Skip to main content

pfx_live/
deform.rs

1use bytemuck::{Pod, Zeroable};
2use pfx_geom::strip::{ROLL_WGSL, Roller};
3use pfx_physics::BEND_WGSL;
4use std::sync::Arc;
5use wgpu::util::DeviceExt;
6
7pub const NONE: u32 = u32::MAX;
8pub const ROLLER: u32 = 1;
9pub const WIND: u32 = 2;
10pub const CURVED_ROLLER: u32 = 3;
11
12pub const DEFORM_WGSL: &str = r#"
13struct DeformerData {
14    kind: vec4u,
15    current_a: vec4f,
16    current_b: vec4f,
17    previous_a: vec4f,
18    previous_b: vec4f,
19}
20struct DeformedVertex {
21    position: vec3f,
22    normal: vec3f,
23    tangent: vec4f,
24    previous_position: vec3f,
25}
26@group(2) @binding(0) var<storage, read> deformers: array<DeformerData>;
27@group(2) @binding(1) var<storage, read> roller_plan: array<vec4f>;
28
29fn plan_sample(s: f32, offset: u32, count: u32) -> vec4f {
30    var low = 0u;
31    var high = count - 1u;
32    let clamped = clamp(s, roller_plan[offset].x, roller_plan[offset + high].x);
33    while (high - low > 1u) {
34        let middle = (low + high) / 2u;
35        if (roller_plan[offset + middle].x < clamped) {
36            low = middle;
37        } else {
38            high = middle;
39        }
40    }
41    let a = roller_plan[offset + low];
42    let b = roller_plan[offset + high];
43    let before = roller_plan[offset + max(low, 1u) - 1u];
44    let after = roller_plan[offset + min(high + 1u, count - 1u)];
45    let span = b.x - a.x;
46    let u = (clamped - a.x) / span;
47    let u2 = u * u;
48    let u3 = u2 * u;
49    let ta = normalize(vec2f(b.y - before.y, b.z - before.z));
50    let tb = normalize(vec2f(after.y - a.y, after.z - a.z));
51    let p = (2.0 * u3 - 3.0 * u2 + 1.0) * a.yz
52        + (u3 - 2.0 * u2 + u) * span * ta
53        + (-2.0 * u3 + 3.0 * u2) * b.yz
54        + (u3 - u2) * span * tb;
55    let derivative = (6.0 * u2 - 6.0 * u) * a.yz / span
56        + (3.0 * u2 - 4.0 * u + 1.0) * ta
57        + (-6.0 * u2 + 6.0 * u) * b.yz / span
58        + (3.0 * u2 - 2.0 * u) * tb;
59    return vec4f(p, normalize(derivative));
60}
61
62fn curved_keep(rest: vec3f, data: DeformerData, previous: bool) -> f32 {
63    let state = select(data.current_a, data.previous_a, previous);
64    return 1.0 - smoothstep(0.0, 3.14159265 * state.y, rest.x - state.x);
65}
66
67fn flattened(rest: vec3f, data: DeformerData, normal: vec3f, tangent: vec3f) -> array<vec3f, 2> {
68    let keep = curved_keep(rest, data, false);
69    if (keep >= 1.0) {
70        return array<vec3f, 2>(normal, tangent);
71    }
72    let side = select(1.0, -1.0, normal.y < 0.0);
73    let n = normalize(mix(vec3f(0.0, side, 0.0), normal, keep));
74    let along = tangent - n * dot(tangent, n);
75    let t = select(normalize(along), vec3f(1.0, 0.0, 0.0), dot(along, along) < 1e-12);
76    return array<vec3f, 2>(n, t);
77}
78
79fn curved_roll(rest: vec3f, data: DeformerData, previous: bool) -> Rolled {
80    let state = select(data.current_a, data.previous_a, previous);
81    if (data.kind.w == 1u) {
82        return roll_place(rest, state.x, state.y, state.z);
83    }
84    let sample = plan_sample(min(rest.x, state.x), data.kind.y, data.kind.z);
85    let direction = sample.zw;
86    let normal = vec2f(-direction.y, direction.x);
87    let rolled = roll_place(rest, state.x, state.y, state.z);
88    var out: Rolled;
89    if (rest.x <= state.x) {
90        out.position = vec3f(sample.xy + normal * rest.y, rest.z);
91        out.tangent = vec3f(direction, 0.0);
92        out.normal = vec3f(normal, 0.0);
93    } else {
94        let travel = rest.x - state.x;
95        let shrink = state.z / 6.28318531;
96        let fade = exp(-travel / state.y);
97        let planar = rolled.position.xy - vec2f(state.x, shrink * (1.0 - fade));
98        let tangent = normalize(rolled.tangent.xy - vec2f(0.0, shrink / state.y * fade));
99        out.position = vec3f(sample.xy + direction * planar.x + normal * planar.y, rest.z);
100        out.tangent = vec3f(direction * tangent.x + normal * tangent.y, 0.0);
101        out.normal = vec3f(direction * -tangent.y + normal * tangent.x, 0.0);
102        out.position += out.normal * rest.y * curved_keep(rest, data, previous);
103    }
104    return out;
105}
106
107fn deform_wind_position(rest: vec3f, data: DeformerData, previous: bool) -> vec3f {
108    let a = select(data.current_a, data.previous_a, previous);
109    let b = select(data.current_b, data.previous_b, previous);
110    return rest + bend(rest, data.kind.y, a.x, a.y, b.xyz, f32(data.kind.z));
111}
112
113fn deform_position(rest: vec3f, deformer_id: u32, previous: bool) -> vec3f {
114    if (deformer_id == 0xffffffffu) {
115        return rest;
116    }
117    let data = deformers[deformer_id];
118    if (data.kind.x == 1u) {
119        let a = select(data.current_a, data.previous_a, previous);
120        return roll_place(rest, a.x, a.y, a.z).position;
121    }
122    if (data.kind.x == 3u) {
123        return curved_roll(rest, data, previous).position;
124    }
125    if (data.kind.x == 2u) {
126        return deform_wind_position(rest, data, previous);
127    }
128    return rest;
129}
130
131fn deform_vertex(rest: vec3f, normal: vec3f, tangent: vec4f, deformer_id: u32) -> DeformedVertex {
132    var out: DeformedVertex;
133    out.position = rest;
134    out.normal = normal;
135    out.tangent = tangent;
136    out.previous_position = rest;
137    if (deformer_id == 0xffffffffu) {
138        return out;
139    }
140    let data = deformers[deformer_id];
141    if (data.kind.x == 1u || data.kind.x == 3u) {
142        var current: Rolled;
143        var n = normal;
144        var t = tangent.xyz;
145        if (data.kind.x == 3u) {
146            current = curved_roll(rest, data, false);
147            if (data.kind.w == 0u) {
148                let flat = flattened(rest, data, normal, tangent.xyz);
149                n = flat[0];
150                t = flat[1];
151            }
152        } else {
153            current = roll_place(rest, data.current_a.x, data.current_a.y, data.current_a.z);
154        }
155        let axis = vec3f(0.0, 0.0, 1.0);
156        out.position = current.position;
157        out.previous_position = deform_position(rest, deformer_id, true);
158        out.normal = normalize(current.tangent * n.x + current.normal * n.y + axis * n.z);
159        out.tangent = vec4f(normalize(current.tangent * t.x + current.normal * t.y + axis * t.z), tangent.w);
160    } else if (data.kind.x == 2u) {
161        out.position = deform_wind_position(rest, data, false);
162        out.previous_position = deform_position(rest, deformer_id, true);
163        let step = 0.001;
164        let dx = (deform_wind_position(rest + vec3f(step, 0.0, 0.0), data, false) - out.position) / step;
165        let dy = (deform_wind_position(rest + vec3f(0.0, step, 0.0), data, false) - out.position) / step;
166        let dz = (deform_wind_position(rest + vec3f(0.0, 0.0, step), data, false) - out.position) / step;
167        let cofactor = mat3x3f(cross(dy, dz), cross(dz, dx), cross(dx, dy));
168        let jacobian = mat3x3f(dx, dy, dz);
169        out.normal = normalize(cofactor * normal);
170        out.tangent = vec4f(normalize(jacobian * tangent.xyz), tangent.w);
171    }
172    return out;
173}
174"#;
175
176pub fn shader_source() -> String {
177    format!("{ROLL_WGSL}\n{BEND_WGSL}\n{DEFORM_WGSL}")
178}
179
180#[derive(Clone, Copy, Debug, PartialEq)]
181pub struct RollerFrame {
182    pub at: f32,
183    pub outer: f32,
184    pub thickness: f32,
185}
186
187#[derive(Clone, Copy, Debug, PartialEq)]
188pub struct WindFrame {
189    pub time: f32,
190    pub strength: f32,
191    pub direction: [f32; 3],
192}
193
194#[derive(Clone, Debug, PartialEq)]
195pub enum Deformer {
196    Roller {
197        current: RollerFrame,
198        previous: RollerFrame,
199    },
200    CurvedRoller {
201        current: RollerFrame,
202        previous: RollerFrame,
203        plan: Arc<[[f32; 4]]>,
204    },
205    Wind {
206        current: WindFrame,
207        previous: WindFrame,
208        object: u32,
209        seed: u32,
210    },
211}
212
213#[derive(Clone, Copy, Debug, PartialEq, Eq)]
214pub struct DeformerId(pub u32);
215
216impl DeformerId {
217    pub const NONE: Self = Self(NONE);
218}
219
220#[repr(C)]
221#[derive(Clone, Copy, Debug, PartialEq, Pod, Zeroable)]
222pub struct PackedDeformer {
223    pub kind: [u32; 4],
224    pub current_a: [f32; 4],
225    pub current_b: [f32; 4],
226    pub previous_a: [f32; 4],
227    pub previous_b: [f32; 4],
228}
229
230impl From<&Deformer> for PackedDeformer {
231    fn from(value: &Deformer) -> Self {
232        match value {
233            Deformer::Roller { current, previous } => Self {
234                kind: [ROLLER, 0, 0, 0],
235                current_a: [current.at, current.outer, current.thickness, 0.0],
236                current_b: [0.0; 4],
237                previous_a: [previous.at, previous.outer, previous.thickness, 0.0],
238                previous_b: [0.0; 4],
239            },
240            Deformer::CurvedRoller {
241                current, previous, ..
242            } => Self {
243                kind: [CURVED_ROLLER, 0, 0, 0],
244                current_a: [current.at, current.outer, current.thickness, 0.0],
245                current_b: [0.0; 4],
246                previous_a: [previous.at, previous.outer, previous.thickness, 0.0],
247                previous_b: [0.0; 4],
248            },
249            Deformer::Wind {
250                current,
251                previous,
252                object,
253                seed,
254            } => Self {
255                kind: [WIND, *object, *seed, 0],
256                current_a: [current.time, current.strength, 0.0, 0.0],
257                current_b: [
258                    current.direction[0],
259                    current.direction[1],
260                    current.direction[2],
261                    0.0,
262                ],
263                previous_a: [previous.time, previous.strength, 0.0, 0.0],
264                previous_b: [
265                    previous.direction[0],
266                    previous.direction[1],
267                    previous.direction[2],
268                    0.0,
269                ],
270            },
271        }
272    }
273}
274
275pub fn pack(deformers: &[Deformer]) -> Result<Vec<PackedDeformer>, String> {
276    pack_with_plans(deformers).map(|(packed, _)| packed)
277}
278
279fn pack_with_plans(deformers: &[Deformer]) -> Result<(Vec<PackedDeformer>, Vec<[f32; 4]>), String> {
280    if deformers.len() >= NONE as usize {
281        return Err("too many deformers".into());
282    }
283    let mut packed = Vec::with_capacity(deformers.len().max(1));
284    let mut plans = Vec::new();
285    for deformer in deformers {
286        match deformer {
287            Deformer::Roller { current, previous }
288            | Deformer::CurvedRoller {
289                current, previous, ..
290            } => {
291                for state in [*current, *previous] {
292                    if !state.at.is_finite()
293                        || !state.outer.is_finite()
294                        || state.outer <= 0.0
295                        || !state.thickness.is_finite()
296                        || state.thickness < 0.0
297                    {
298                        return Err("roller parameters must be finite with positive radius and nonnegative thickness".into());
299                    }
300                }
301            }
302            Deformer::Wind {
303                current, previous, ..
304            } => {
305                for state in [current, previous] {
306                    if !state.time.is_finite()
307                        || !state.strength.is_finite()
308                        || !state.direction.iter().all(|value| value.is_finite())
309                    {
310                        return Err("wind parameters must be finite".into());
311                    }
312                }
313            }
314        }
315        let mut item = PackedDeformer::from(deformer);
316        if let Deformer::CurvedRoller { plan, .. } = deformer {
317            if plan.len() < 2 || plan.len() > u32::MAX as usize {
318                return Err("roller plan needs at least two samples".into());
319            }
320            for (index, point) in plan.iter().enumerate() {
321                if !point.iter().all(|value| value.is_finite())
322                    || (index > 0 && point[0] <= plan[index - 1][0])
323                    || (index > 0
324                        && (point[1] - plan[index - 1][1]).hypot(point[2] - plan[index - 1][2])
325                            <= 1e-9)
326                {
327                    return Err("roller plan must have finite, increasing arc lengths and distinct positions".into());
328                }
329            }
330            if plans.len() > u32::MAX as usize - plan.len() {
331                return Err("roller plan buffer is too large".into());
332            }
333            item.kind[1] = plans.len() as u32;
334            item.kind[2] = plan.len() as u32;
335            item.kind[3] = u32::from(
336                plan.iter()
337                    .all(|point| point[1] == point[0] && point[2] == 0.0),
338            );
339            plans.extend_from_slice(plan);
340        }
341        packed.push(item);
342    }
343    if packed.is_empty() {
344        packed.push(PackedDeformer::zeroed());
345    }
346    if plans.is_empty() {
347        plans.push([0.0; 4]);
348    }
349    Ok((packed, plans))
350}
351
352pub struct DeformerBuffer {
353    pub layout: wgpu::BindGroupLayout,
354    pub group: wgpu::BindGroup,
355    buffer: wgpu::Buffer,
356    capacity: usize,
357    plan_buffer: wgpu::Buffer,
358    plan_capacity: usize,
359    pub len: usize,
360}
361
362impl DeformerBuffer {
363    pub fn new(device: &wgpu::Device) -> Self {
364        let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
365            label: Some("live deformers"),
366            entries: &deformer_entries(),
367        });
368        let buffer = device.create_buffer(&wgpu::BufferDescriptor {
369            label: Some("live deformer parameters"),
370            size: std::mem::size_of::<PackedDeformer>() as u64,
371            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
372            mapped_at_creation: false,
373        });
374        let plan_buffer = device.create_buffer(&wgpu::BufferDescriptor {
375            label: Some("live roller plan"),
376            size: 16,
377            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
378            mapped_at_creation: false,
379        });
380        let group = Self::group(device, &layout, &buffer, &plan_buffer);
381        Self {
382            layout,
383            group,
384            buffer,
385            capacity: 1,
386            plan_buffer,
387            plan_capacity: 1,
388            len: 0,
389        }
390    }
391
392    pub fn buffer(&self) -> &wgpu::Buffer {
393        &self.buffer
394    }
395
396    pub fn plan_buffer(&self) -> &wgpu::Buffer {
397        &self.plan_buffer
398    }
399
400    fn group(
401        device: &wgpu::Device,
402        layout: &wgpu::BindGroupLayout,
403        buffer: &wgpu::Buffer,
404        plan_buffer: &wgpu::Buffer,
405    ) -> wgpu::BindGroup {
406        device.create_bind_group(&wgpu::BindGroupDescriptor {
407            label: Some("live deformers"),
408            layout,
409            entries: &[
410                wgpu::BindGroupEntry {
411                    binding: 0,
412                    resource: buffer.as_entire_binding(),
413                },
414                wgpu::BindGroupEntry {
415                    binding: 1,
416                    resource: plan_buffer.as_entire_binding(),
417                },
418            ],
419        })
420    }
421
422    pub fn upload(
423        &mut self,
424        device: &wgpu::Device,
425        queue: &wgpu::Queue,
426        deformers: &[Deformer],
427    ) -> Result<(), String> {
428        let (packed, plans) = pack_with_plans(deformers)?;
429        let mut regroup = false;
430        if packed.len() > self.capacity {
431            self.buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
432                label: Some("live deformer parameters"),
433                contents: bytemuck::cast_slice(&packed),
434                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
435            });
436            self.capacity = packed.len();
437            regroup = true;
438        } else {
439            queue.write_buffer(&self.buffer, 0, bytemuck::cast_slice(&packed));
440        }
441        if plans.len() > self.plan_capacity {
442            self.plan_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
443                label: Some("live roller plan"),
444                contents: bytemuck::cast_slice(&plans),
445                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
446            });
447            self.plan_capacity = plans.len();
448            regroup = true;
449        } else {
450            queue.write_buffer(&self.plan_buffer, 0, bytemuck::cast_slice(&plans));
451        }
452        if regroup {
453            self.group = Self::group(device, &self.layout, &self.buffer, &self.plan_buffer);
454        }
455        self.len = deformers.len();
456        Ok(())
457    }
458}
459
460#[derive(Clone, Copy, Debug, PartialEq)]
461pub struct DeformedRollerVertex {
462    pub position: [f32; 3],
463    pub normal: [f32; 3],
464    pub tangent: [f32; 4],
465    pub previous_position: [f32; 3],
466    pub uv: [f32; 2],
467}
468
469pub fn roller_vertex(
470    rest: [f32; 3],
471    normal: [f32; 3],
472    tangent: [f32; 4],
473    uv: [f32; 2],
474    current: RollerFrame,
475    previous: RollerFrame,
476) -> DeformedRollerVertex {
477    let make = |state: RollerFrame| {
478        Roller {
479            at: state.at,
480            core: state.outer,
481            thickness: state.thickness,
482        }
483        .place(rest, state.outer)
484    };
485    let now = make(current);
486    let before = make(previous);
487    let rotate = |v: [f32; 3]| {
488        [
489            now.tangent[0] * v[0] + now.normal[0] * v[1],
490            now.tangent[1] * v[0] + now.normal[1] * v[1],
491            v[2],
492        ]
493    };
494    let n = rotate(normal);
495    let t = rotate([tangent[0], tangent[1], tangent[2]]);
496    DeformedRollerVertex {
497        position: now.position,
498        normal: n,
499        tangent: [t[0], t[1], t[2], tangent[3]],
500        previous_position: before.position,
501        uv,
502    }
503}
504
505pub fn plan_sample(plan: &[[f32; 4]], s: f32) -> ([f32; 2], [f32; 2]) {
506    let s = s.clamp(plan[0][0], plan[plan.len() - 1][0]);
507    let high = plan
508        .partition_point(|point| point[0] < s)
509        .clamp(1, plan.len() - 1);
510    let low = high - 1;
511    let a = plan[low];
512    let b = plan[high];
513    let before = plan[low.saturating_sub(1)];
514    let after = plan[(high + 1).min(plan.len() - 1)];
515    let span = b[0] - a[0];
516    let u = (s - a[0]) / span;
517    let u2 = u * u;
518    let u3 = u2 * u;
519    let unit = |x: f32, y: f32| {
520        let length = x.hypot(y);
521        [x / length, y / length]
522    };
523    let ta = unit(b[1] - before[1], b[2] - before[2]);
524    let tb = unit(after[1] - a[1], after[2] - a[2]);
525    let position = std::array::from_fn(|i| {
526        let k = i + 1;
527        (2.0 * u3 - 3.0 * u2 + 1.0) * a[k]
528            + (u3 - 2.0 * u2 + u) * span * ta[i]
529            + (-2.0 * u3 + 3.0 * u2) * b[k]
530            + (u3 - u2) * span * tb[i]
531    });
532    let derivative: [f32; 2] = std::array::from_fn(|i| {
533        let k = i + 1;
534        (6.0 * u2 - 6.0 * u) * a[k] / span
535            + (3.0 * u2 - 4.0 * u + 1.0) * ta[i]
536            + (-6.0 * u2 + 6.0 * u) * b[k] / span
537            + (3.0 * u2 - 2.0 * u) * tb[i]
538    });
539    (position, unit(derivative[0], derivative[1]))
540}
541
542fn smoothstep(low: f32, high: f32, x: f32) -> f32 {
543    let t = ((x - low) / (high - low)).clamp(0.0, 1.0);
544    t * t * (3.0 - 2.0 * t)
545}
546
547fn flattened(keep: f32, normal: [f32; 3], tangent: [f32; 3]) -> ([f32; 3], [f32; 3]) {
548    if keep >= 1.0 {
549        return (normal, tangent);
550    }
551    let side = if normal[1] < 0.0 { -1.0 } else { 1.0 };
552    let mixed = [
553        normal[0] * keep,
554        side + (normal[1] - side) * keep,
555        normal[2] * keep,
556    ];
557    let length = mixed.iter().map(|v| v * v).sum::<f32>().sqrt();
558    let n = mixed.map(|v| v / length);
559    let along_n = tangent.iter().zip(n).map(|(a, b)| a * b).sum::<f32>();
560    let along: [f32; 3] = std::array::from_fn(|i| tangent[i] - n[i] * along_n);
561    let squared = along.iter().map(|v| v * v).sum::<f32>();
562    let t = if squared < 1e-12 {
563        [1.0, 0.0, 0.0]
564    } else {
565        along.map(|v| v / squared.sqrt())
566    };
567    (n, t)
568}
569
570pub fn curved_roller_vertex(
571    rest: [f32; 3],
572    normal: [f32; 3],
573    tangent: [f32; 4],
574    uv: [f32; 2],
575    current: RollerFrame,
576    previous: RollerFrame,
577    plan: &[[f32; 4]],
578) -> DeformedRollerVertex {
579    if plan
580        .iter()
581        .all(|point| point[1] == point[0] && point[2] == 0.0)
582    {
583        return roller_vertex(rest, normal, tangent, uv, current, previous);
584    }
585    let keep = |state: RollerFrame| {
586        1.0 - smoothstep(0.0, std::f32::consts::PI * state.outer, rest[0] - state.at)
587    };
588    let place = |state: RollerFrame| {
589        let (anchor, direction) = plan_sample(plan, rest[0].min(state.at));
590        let side = [-direction[1], direction[0]];
591        let roll = Roller {
592            at: state.at,
593            core: state.outer,
594            thickness: state.thickness,
595        }
596        .place(rest, state.outer);
597        if rest[0] <= state.at {
598            (
599                [
600                    anchor[0] + side[0] * rest[1],
601                    anchor[1] + side[1] * rest[1],
602                    rest[2],
603                ],
604                [direction[0], direction[1]],
605                side,
606            )
607        } else {
608            let dx = roll.position[0] - state.at;
609            let travel = rest[0] - state.at;
610            let shrink = state.thickness / std::f32::consts::TAU;
611            let fade = (-travel / state.outer).exp();
612            let dy = roll.position[1] - shrink * (1.0 - fade);
613            let raw = [
614                roll.tangent[0],
615                roll.tangent[1] - shrink / state.outer * fade,
616            ];
617            let length = raw[0].hypot(raw[1]);
618            let local_tangent = [raw[0] / length, raw[1] / length];
619            let tangent = [
620                direction[0] * local_tangent[0] + side[0] * local_tangent[1],
621                direction[1] * local_tangent[0] + side[1] * local_tangent[1],
622            ];
623            let normal = [-tangent[1], tangent[0]];
624            let lift = rest[1] * keep(state);
625            (
626                [
627                    anchor[0] + direction[0] * dx + side[0] * dy + normal[0] * lift,
628                    anchor[1] + direction[1] * dx + side[1] * dy + normal[1] * lift,
629                    rest[2],
630                ],
631                tangent,
632                normal,
633            )
634        }
635    };
636    let now = place(current);
637    let before = place(previous);
638    let rotate = |v: [f32; 3]| {
639        [
640            now.1[0] * v[0] + now.2[0] * v[1],
641            now.1[1] * v[0] + now.2[1] * v[1],
642            v[2],
643        ]
644    };
645    let (flat_normal, flat_tangent) =
646        flattened(keep(current), normal, [tangent[0], tangent[1], tangent[2]]);
647    let n = rotate(flat_normal);
648    let t = rotate(flat_tangent);
649    DeformedRollerVertex {
650        position: now.0,
651        normal: n,
652        tangent: [t[0], t[1], t[2], tangent[3]],
653        previous_position: before.0,
654        uv,
655    }
656}
657
658pub fn deformer_entries() -> Vec<wgpu::BindGroupLayoutEntry> {
659    vec![
660        wgpu::BindGroupLayoutEntry {
661            binding: 0,
662            visibility: wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::COMPUTE,
663            ty: wgpu::BindingType::Buffer {
664                ty: wgpu::BufferBindingType::Storage { read_only: true },
665                has_dynamic_offset: false,
666                min_binding_size: None,
667            },
668            count: None,
669        },
670        wgpu::BindGroupLayoutEntry {
671            binding: 1,
672            visibility: wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::COMPUTE,
673            ty: wgpu::BindingType::Buffer {
674                ty: wgpu::BufferBindingType::Storage { read_only: true },
675                has_dynamic_offset: false,
676                min_binding_size: None,
677            },
678            count: None,
679        },
680    ]
681}
682
683#[cfg(test)]
684mod tests {
685    use super::*;
686    use pfx_geom::strip::Strip;
687
688    #[test]
689    fn parameter_packing() {
690        assert_eq!(std::mem::size_of::<PackedDeformer>(), 80);
691        assert_eq!(DeformerId::NONE.0, NONE);
692        let roller = Deformer::Roller {
693            current: RollerFrame {
694                at: 0.3,
695                outer: 0.2,
696                thickness: 0.01,
697            },
698            previous: RollerFrame {
699                at: 0.4,
700                outer: 0.22,
701                thickness: 0.01,
702            },
703        };
704        let wind = Deformer::Wind {
705            current: WindFrame {
706                time: 1.5,
707                strength: 2.0,
708                direction: [1.0, 0.0, 0.0],
709            },
710            previous: WindFrame {
711                time: 1.0,
712                strength: 1.5,
713                direction: [0.0, 0.0, 1.0],
714            },
715            object: 42,
716            seed: 7,
717        };
718        let packed = pack(&[roller, wind]).unwrap();
719        assert_eq!(packed[0].kind, [ROLLER, 0, 0, 0]);
720        assert_eq!(packed[0].current_a, [0.3, 0.2, 0.01, 0.0]);
721        assert_eq!(packed[0].previous_a, [0.4, 0.22, 0.01, 0.0]);
722        assert_eq!(packed[1].kind, [WIND, 42, 7, 0]);
723        assert_eq!(packed[1].current_b, [1.0, 0.0, 0.0, 0.0]);
724        assert_eq!(packed[1].previous_b, [0.0, 0.0, 1.0, 0.0]);
725        assert_eq!(bytemuck::bytes_of(&packed[0]).len(), 80);
726        assert!(
727            pack(&[Deformer::Roller {
728                current: RollerFrame {
729                    at: 0.3,
730                    outer: 0.0,
731                    thickness: 0.01
732                },
733                previous: RollerFrame {
734                    at: 0.4,
735                    outer: 0.2,
736                    thickness: 0.01
737                },
738            }])
739            .is_err()
740        );
741    }
742
743    #[test]
744    fn roller_matches_geom_and_keeps_uv() {
745        let current = RollerFrame {
746            at: 0.4,
747            outer: 0.17,
748            thickness: 0.008,
749        };
750        let previous = RollerFrame {
751            at: 0.48,
752            outer: 0.18,
753            thickness: 0.008,
754        };
755        let roller = Roller {
756            at: current.at,
757            core: current.outer,
758            thickness: current.thickness,
759        };
760        let before = Roller {
761            at: previous.at,
762            core: previous.outer,
763            thickness: previous.thickness,
764        };
765        for rest in [
766            [0.1, 0.0, 0.2],
767            [0.4, 0.0, 0.5],
768            [0.51, 0.0, 0.5],
769            [0.8, 0.0, 0.8],
770        ] {
771            let uv = [rest[0], rest[2]];
772            let got = roller_vertex(
773                rest,
774                [0.0, 1.0, 0.0],
775                [1.0, 0.0, 0.0, 1.0],
776                uv,
777                current,
778                previous,
779            );
780            let reference = roller.place(rest, current.outer);
781            let previous_reference = before.place(rest, previous.outer);
782            assert_eq!(got.position, reference.position);
783            assert_eq!(got.normal, reference.normal);
784            assert_eq!(
785                got.tangent,
786                [
787                    reference.tangent[0],
788                    reference.tangent[1],
789                    reference.tangent[2],
790                    1.0
791                ]
792            );
793            assert_eq!(got.previous_position, previous_reference.position);
794            assert_eq!(got.uv, uv);
795            if rest[0] > current.at {
796                assert_ne!(got.position, got.previous_position);
797            }
798        }
799    }
800
801    #[test]
802    fn straight_plan_matches_flat_roller() {
803        let plan = [
804            [0.0, 0.0, 0.0, 0.0],
805            [0.5, 0.5, 0.0, 0.0],
806            [1.0, 1.0, 0.0, 0.0],
807        ];
808        let current = RollerFrame {
809            at: 0.4,
810            outer: 0.17,
811            thickness: 0.008,
812        };
813        let previous = RollerFrame {
814            at: 0.48,
815            outer: 0.18,
816            thickness: 0.008,
817        };
818        let packed = pack(&[Deformer::CurvedRoller {
819            current,
820            previous,
821            plan: plan.to_vec().into(),
822        }])
823        .unwrap();
824        assert_eq!(packed[0].kind[3], 1);
825        for s in [0.0, 0.25, 0.4, 0.5, 0.8, 1.0] {
826            let rest = [s, 0.0, 0.2];
827            let flat = roller_vertex(
828                rest,
829                [0.0, 1.0, 0.0],
830                [1.0, 0.0, 0.0, 1.0],
831                [s, 0.2],
832                current,
833                previous,
834            );
835            let curved = curved_roller_vertex(
836                rest,
837                [0.0, 1.0, 0.0],
838                [1.0, 0.0, 0.0, 1.0],
839                [s, 0.2],
840                current,
841                previous,
842                &plan,
843            );
844            for (a, b) in flat.position.iter().zip(curved.position) {
845                assert!((a - b).abs() < 1e-6);
846            }
847            for (a, b) in flat.previous_position.iter().zip(curved.previous_position) {
848                assert!((a - b).abs() < 1e-6);
849            }
850            assert_eq!(flat.normal, curved.normal);
851            assert_eq!(flat.tangent, curved.tangent);
852        }
853    }
854
855    #[test]
856    fn curved_plan_packs_arc_length_samples() {
857        let plan: Arc<[[f32; 4]]> = vec![[0.0, 0.0, 0.0, 0.0], [0.5, 0.49, 0.05, 0.0]].into();
858        let state = RollerFrame {
859            at: 0.4,
860            outer: 0.02,
861            thickness: 0.002,
862        };
863        let roller = Deformer::CurvedRoller {
864            current: state,
865            previous: state,
866            plan: plan.clone(),
867        };
868        let (packed, samples) = pack_with_plans(&[roller.clone(), roller]).unwrap();
869        assert_eq!(packed[0].kind, [CURVED_ROLLER, 0, 2, 0]);
870        assert_eq!(packed[1].kind, [CURVED_ROLLER, 2, 2, 0]);
871        assert_eq!(samples, [&plan[..], &plan[..]].concat());
872        let bad: Arc<[[f32; 4]]> = vec![[0.0, 0.0, 0.0, 0.0], [0.0, 1.0, 0.0, 0.0]].into();
873        assert!(
874            pack(&[Deformer::CurvedRoller {
875                current: state,
876                previous: state,
877                plan: bad
878            }])
879            .is_err()
880        );
881    }
882
883    #[test]
884    fn curved_plan_samples_and_join_are_continuous() {
885        let plan = [
886            [0.0, 0.0, 0.0, 0.0],
887            [0.2, 0.19, 0.05, 0.0],
888            [0.4, 0.38, 0.02, 0.0],
889            [0.6, 0.58, 0.0, 0.0],
890        ];
891        for knot in plan.iter().skip(1).take(2) {
892            let left = plan_sample(&plan, knot[0] - 1e-5);
893            let middle = plan_sample(&plan, knot[0]);
894            let right = plan_sample(&plan, knot[0] + 1e-5);
895            assert!((middle.0[0] - knot[1]).abs() < 1e-6);
896            assert!((middle.0[1] - knot[2]).abs() < 1e-6);
897            for i in 0..2 {
898                assert!((left.1[i] - right.1[i]).abs() < 1e-3);
899            }
900        }
901        let state = RollerFrame {
902            at: 0.3,
903            outer: 0.02,
904            thickness: 0.002,
905        };
906        let get = |s| {
907            curved_roller_vertex(
908                [s, 0.0, 0.1],
909                [0.0, 1.0, 0.0],
910                [1.0, 0.0, 0.0, 1.0],
911                [s, 0.1],
912                state,
913                state,
914                &plan,
915            )
916        };
917        let left = get(state.at - 1e-6);
918        let right = get(state.at + 1e-6);
919        for i in 0..3 {
920            assert!((left.position[i] - right.position[i]).abs() < 1e-5);
921            assert!((left.normal[i] - right.normal[i]).abs() < 1e-3);
922            assert!((left.tangent[i] - right.tangent[i]).abs() < 1e-3);
923        }
924    }
925
926    #[test]
927    fn moving_fold_has_previous_position_and_attached_glyphs() {
928        let plan = [
929            [0.0, 0.0, 0.0, 0.0],
930            [0.3, 0.29, 0.04, 0.0],
931            [0.6, 0.58, 0.0, 0.0],
932        ];
933        let current = RollerFrame {
934            at: 0.32,
935            outer: 0.018,
936            thickness: 0.002,
937        };
938        let previous = RollerFrame {
939            at: 0.42,
940            outer: 0.016,
941            thickness: 0.002,
942        };
943        let glyph = [0.36, 0.0, 0.14];
944        let ink = curved_roller_vertex(
945            glyph,
946            [0.0, 1.0, 0.0],
947            [1.0, 0.0, 0.0, 1.0],
948            [glyph[0], glyph[2]],
949            current,
950            previous,
951            &plan,
952        );
953        let offset = ((ink.position[0] - ink.previous_position[0]).powi(2)
954            + (ink.position[1] - ink.previous_position[1]).powi(2))
955        .sqrt();
956        assert!(offset > 0.001);
957        assert_eq!(ink.uv, [glyph[0], glyph[2]]);
958        let lifted = |s: f32, lift: f32| {
959            curved_roller_vertex(
960                [s, lift, glyph[2]],
961                [0.0, 1.0, 0.0],
962                [1.0, 0.0, 0.0, 1.0],
963                [s, glyph[2]],
964                current,
965                previous,
966                &plan,
967            )
968        };
969        let open = lifted(0.2, 0.0);
970        let raised = lifted(0.2, 0.004);
971        for i in 0..3 {
972            assert!((raised.position[i] - open.position[i] - open.normal[i] * 0.004).abs() < 1e-6);
973        }
974        let wound = current.at + std::f32::consts::PI * current.outer;
975        for s in [wound, wound + 0.05] {
976            let flat = lifted(s, 0.0);
977            let raised = lifted(s, 0.004);
978            assert_eq!(raised.position, flat.position);
979        }
980        let joined = |lift: f32| {
981            let left = lifted(current.at - 1e-6, lift);
982            let right = lifted(current.at + 1e-6, lift);
983            (0..3).all(|i| (left.position[i] - right.position[i]).abs() < 1e-5)
984        };
985        assert!(joined(0.004));
986        let tilted = curved_roller_vertex(
987            [wound + 0.01, 0.004, glyph[2]],
988            [0.6, 0.8, 0.0],
989            [0.8, -0.6, 0.0, 1.0],
990            ink.uv,
991            current,
992            previous,
993            &plan,
994        );
995        let flat = lifted(wound + 0.01, 0.0);
996        for i in 0..3 {
997            assert!((tilted.normal[i] - flat.normal[i]).abs() < 1e-5);
998            assert!((tilted.tangent[i] - flat.tangent[i]).abs() < 1e-5);
999        }
1000    }
1001
1002    #[test]
1003    #[ignore = "needs a GPU; run under pgpu"]
1004    fn curved_glyph_shader_matches_cpu_mirror() {
1005        let gpu = pollster::block_on(pfx_gpu::Gpu::headless()).unwrap();
1006        let device = &gpu.device;
1007        let plan: Arc<[[f32; 4]]> = vec![
1008            [0.0, 0.0, 0.0, 0.0],
1009            [0.2, 0.19, 0.04, 0.0],
1010            [0.4, 0.38, 0.02, 0.0],
1011            [0.6, 0.58, 0.0, 0.0],
1012        ]
1013        .into();
1014        let current = RollerFrame {
1015            at: 0.32,
1016            outer: 0.018,
1017            thickness: 0.002,
1018        };
1019        let previous = RollerFrame {
1020            at: 0.42,
1021            outer: 0.016,
1022            thickness: 0.002,
1023        };
1024        let mut deformers = DeformerBuffer::new(device);
1025        deformers
1026            .upload(
1027                device,
1028                &gpu.queue,
1029                &[Deformer::CurvedRoller {
1030                    current,
1031                    previous,
1032                    plan: plan.clone(),
1033                }],
1034            )
1035            .unwrap();
1036        let output = device.create_buffer(&wgpu::BufferDescriptor {
1037            label: Some("curved glyph positions"),
1038            size: 192,
1039            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
1040            mapped_at_creation: false,
1041        });
1042        let readback = device.create_buffer(&wgpu::BufferDescriptor {
1043            label: Some("curved glyph readback"),
1044            size: 192,
1045            usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
1046            mapped_at_creation: false,
1047        });
1048        let result_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
1049            label: None,
1050            entries: &[wgpu::BindGroupLayoutEntry {
1051                binding: 0,
1052                visibility: wgpu::ShaderStages::COMPUTE,
1053                ty: wgpu::BindingType::Buffer {
1054                    ty: wgpu::BufferBindingType::Storage { read_only: false },
1055                    has_dynamic_offset: false,
1056                    min_binding_size: None,
1057                },
1058                count: None,
1059            }],
1060        });
1061        let result_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
1062            label: None,
1063            layout: &result_layout,
1064            entries: &[wgpu::BindGroupEntry {
1065                binding: 0,
1066                resource: output.as_entire_binding(),
1067            }],
1068        });
1069        let empty = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
1070            label: None,
1071            entries: &[],
1072        });
1073        let empty_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
1074            label: None,
1075            layout: &empty,
1076            entries: &[],
1077        });
1078        let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
1079            label: None,
1080            source: wgpu::ShaderSource::Wgsl(
1081                format!(
1082                    "{}\n{}",
1083                    shader_source(),
1084                    r#"
1085@group(0) @binding(0) var<storage, read_write> result: array<vec4f>;
1086@compute @workgroup_size(3)
1087fn main(@builtin(local_invocation_index) index: u32) {
1088    let rest = vec3f(0.30 + f32(index) * 0.04, 0.004, 0.14);
1089    let shape = deform_vertex(rest, normalize(vec3f(0.3, 1.0, 0.2)), vec4f(normalize(vec3f(1.0, -0.3, 0.0)), 1.0), 0u);
1090    result[index * 4u] = vec4f(shape.position, 1.0);
1091    result[index * 4u + 1u] = vec4f(shape.previous_position, 1.0);
1092    result[index * 4u + 2u] = vec4f(shape.normal, 0.0);
1093    result[index * 4u + 3u] = shape.tangent;
1094}
1095"#
1096                )
1097                .into(),
1098            ),
1099        });
1100        let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
1101            label: None,
1102            bind_group_layouts: &[Some(&result_layout), Some(&empty), Some(&deformers.layout)],
1103            immediate_size: 0,
1104        });
1105        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
1106            label: None,
1107            layout: Some(&layout),
1108            module: &module,
1109            entry_point: Some("main"),
1110            compilation_options: Default::default(),
1111            cache: None,
1112        });
1113        let mut encoder = device.create_command_encoder(&Default::default());
1114        {
1115            let mut pass = encoder.begin_compute_pass(&Default::default());
1116            pass.set_pipeline(&pipeline);
1117            pass.set_bind_group(0, &result_group, &[]);
1118            pass.set_bind_group(1, &empty_group, &[]);
1119            pass.set_bind_group(2, &deformers.group, &[]);
1120            pass.dispatch_workgroups(1, 1, 1);
1121        }
1122        encoder.copy_buffer_to_buffer(&output, 0, &readback, 0, 192);
1123        gpu.queue.submit(Some(encoder.finish()));
1124        let (send, receive) = std::sync::mpsc::channel();
1125        let slice = readback.slice(..);
1126        slice.map_async(wgpu::MapMode::Read, move |result| {
1127            let _ = send.send(result);
1128        });
1129        device.poll(wgpu::PollType::wait_indefinitely()).unwrap();
1130        receive.recv().unwrap().unwrap();
1131        let mapped = slice.get_mapped_range().expect("a mapped range");
1132        let values: &[[f32; 4]] = bytemuck::cast_slice(&mapped[..]);
1133        let unit = |v: [f32; 3]| {
1134            let length = v.iter().map(|x| x * x).sum::<f32>().sqrt();
1135            v.map(|x| x / length)
1136        };
1137        let normal = unit([0.3, 1.0, 0.2]);
1138        let tangent = unit([1.0, -0.3, 0.0]);
1139        for index in 0..3 {
1140            let s = 0.30 + index as f32 * 0.04;
1141            let mirror = curved_roller_vertex(
1142                [s, 0.004, 0.14],
1143                normal,
1144                [tangent[0], tangent[1], tangent[2], 1.0],
1145                [s, 0.14],
1146                current,
1147                previous,
1148                &plan,
1149            );
1150            let close = |actual: &[f32], expected: &[f32], tolerance: f32| {
1151                actual
1152                    .iter()
1153                    .zip(expected)
1154                    .all(|(a, e)| (a - e).abs() < tolerance)
1155            };
1156            assert!(close(&values[index * 4][..3], &mirror.position, 2e-5));
1157            assert!(close(
1158                &values[index * 4 + 1][..3],
1159                &mirror.previous_position,
1160                2e-5
1161            ));
1162            assert!(close(&values[index * 4 + 2][..3], &mirror.normal, 1e-4));
1163            assert!(close(&values[index * 4 + 3], &mirror.tangent, 1e-4));
1164        }
1165        drop(mapped);
1166        readback.unmap();
1167    }
1168
1169    #[test]
1170    fn shader_has_shared_deformation() {
1171        let source = format!(
1172            "{}\n@vertex fn main(@builtin(vertex_index) index: u32) -> @builtin(position) vec4f {{ let d = deform_vertex(vec3f(f32(index), 0.0, 0.0), vec3f(0.0, 1.0, 0.0), vec4f(1.0, 0.0, 0.0, 1.0), 0u); return vec4f(d.position, 1.0); }}",
1173            shader_source()
1174        );
1175        naga::front::wgsl::parse_str(&source).unwrap();
1176    }
1177
1178    #[repr(C)]
1179    #[derive(Clone, Copy, Pod, Zeroable)]
1180    struct TestVertex {
1181        rest: [f32; 4],
1182        normal: [f32; 4],
1183        tangent: [f32; 4],
1184        uv: [f32; 4],
1185    }
1186
1187    const TEST_SHADER: &str = r#"
1188struct Input {
1189    @location(0) rest: vec4f,
1190    @location(1) normal: vec4f,
1191    @location(2) tangent: vec4f,
1192    @location(3) uv: vec4f,
1193}
1194struct Output {
1195    @builtin(position) clip: vec4f,
1196    @location(0) uv: vec2f,
1197    @location(1) height: f32,
1198    @location(2) shadow_difference: f32,
1199}
1200@vertex fn vs_main(input: Input) -> Output {
1201    let v = deform_vertex(input.rest.xyz, input.normal.xyz, input.tangent, 0u);
1202    let shadow_position = deform_position(input.rest.xyz, 0u, false);
1203    var out: Output;
1204    out.clip = vec4f(v.position.x * 2.0 - 1.0, v.position.z * 2.0 - 1.0, 0.8 - v.position.y, 1.0);
1205    out.uv = input.uv.xy;
1206    out.height = v.position.y;
1207    out.shadow_difference = distance(v.position, shadow_position);
1208    return out;
1209}
1210@fragment fn fs_main(input: Output) -> @location(0) vec4f {
1211    if (input.shadow_difference > 0.00001) {
1212        return vec4f(1.0, 0.0, 1.0, 1.0);
1213    }
1214    let ink = abs(fract(input.uv.x * 10.0) - 0.5) < 0.12 && input.uv.y > 0.25 && input.uv.y < 0.75;
1215    let colour = select(vec3f(0.95, 0.9, 0.8), vec3f(0.0), ink);
1216    return vec4f(colour, clamp(input.height * 2.0, 0.0, 1.0));
1217}
1218"#;
1219
1220    #[test]
1221    #[ignore = "needs a GPU; run under pgpu"]
1222    fn rolled_strip_keeps_ink_on_raised_paper() {
1223        let gpu = pollster::block_on(pfx_gpu::Gpu::headless()).unwrap();
1224        let device = &gpu.device;
1225        let strip = Strip {
1226            length: 1.0,
1227            width: 1.0,
1228        };
1229        let mesh = strip.mesh(0.17, 0.0003, None, 12);
1230        let vertices: Vec<TestVertex> = mesh
1231            .positions
1232            .iter()
1233            .enumerate()
1234            .map(|(index, rest)| TestVertex {
1235                rest: [rest[0], rest[1], rest[2], 1.0],
1236                normal: [0.0, 1.0, 0.0, 0.0],
1237                tangent: [1.0, 0.0, 0.0, 1.0],
1238                uv: [mesh.uvs[index][0], mesh.uvs[index][1], 0.0, 0.0],
1239            })
1240            .collect();
1241        let vertex_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
1242            label: Some("deform test strip"),
1243            contents: bytemuck::cast_slice(&vertices),
1244            usage: wgpu::BufferUsages::VERTEX,
1245        });
1246        let index_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
1247            label: Some("deform test indices"),
1248            contents: bytemuck::cast_slice(&mesh.indices),
1249            usage: wgpu::BufferUsages::INDEX,
1250        });
1251        let mut deformers = DeformerBuffer::new(device);
1252        let empty = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
1253            label: None,
1254            entries: &[],
1255        });
1256        let empty_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
1257            label: None,
1258            layout: &empty,
1259            entries: &[],
1260        });
1261        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
1262            label: Some("deform test shader"),
1263            source: wgpu::ShaderSource::Wgsl(
1264                format!("{}\n{}", shader_source(), TEST_SHADER).into(),
1265            ),
1266        });
1267        let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
1268            label: None,
1269            bind_group_layouts: &[Some(&empty), Some(&empty), Some(&deformers.layout)],
1270            immediate_size: 0,
1271        });
1272        let pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
1273            label: Some("deform test pipeline"),
1274            layout: Some(&layout),
1275            vertex: wgpu::VertexState {
1276                module: &shader,
1277                entry_point: Some("vs_main"),
1278                compilation_options: Default::default(),
1279                buffers: &[Some(wgpu::VertexBufferLayout {
1280                    array_stride: std::mem::size_of::<TestVertex>() as u64,
1281                    step_mode: wgpu::VertexStepMode::Vertex,
1282                    attributes: &wgpu::vertex_attr_array![0 => Float32x4, 1 => Float32x4, 2 => Float32x4, 3 => Float32x4],
1283                })],
1284            },
1285            primitive: wgpu::PrimitiveState { cull_mode: None, ..Default::default() },
1286            depth_stencil: Some(wgpu::DepthStencilState {
1287                format: wgpu::TextureFormat::Depth32Float,
1288                depth_write_enabled: Some(true),
1289                depth_compare: Some(wgpu::CompareFunction::Less),
1290                stencil: Default::default(),
1291                bias: Default::default(),
1292            }),
1293            multisample: Default::default(),
1294            fragment: Some(wgpu::FragmentState {
1295                module: &shader,
1296                entry_point: Some("fs_main"),
1297                compilation_options: Default::default(),
1298                targets: &[Some(wgpu::ColorTargetState {
1299                    format: wgpu::TextureFormat::Rgba8UnormSrgb,
1300                    blend: None,
1301                    write_mask: wgpu::ColorWrites::ALL,
1302                })],
1303            }),
1304            multiview_mask: None,
1305            cache: None,
1306        });
1307        let target = gpu
1308            .offscreen(256, 256, wgpu::TextureFormat::Rgba8UnormSrgb)
1309            .unwrap();
1310        let depth = device.create_texture(&wgpu::TextureDescriptor {
1311            label: None,
1312            size: wgpu::Extent3d {
1313                width: 256,
1314                height: 256,
1315                depth_or_array_layers: 1,
1316            },
1317            mip_level_count: 1,
1318            sample_count: 1,
1319            dimension: wgpu::TextureDimension::D2,
1320            format: wgpu::TextureFormat::Depth32Float,
1321            usage: wgpu::TextureUsages::RENDER_ATTACHMENT,
1322            view_formats: &[],
1323        });
1324        let depth_view = depth.create_view(&Default::default());
1325        for at in [0.2, 0.4, 0.6] {
1326            let current = RollerFrame {
1327                at,
1328                outer: 0.17,
1329                thickness: 0.004,
1330            };
1331            deformers
1332                .upload(
1333                    device,
1334                    &gpu.queue,
1335                    &[Deformer::Roller {
1336                        current,
1337                        previous: current,
1338                    }],
1339                )
1340                .unwrap();
1341            let mut encoder = device.create_command_encoder(&Default::default());
1342            {
1343                let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
1344                    multiview_mask: None,
1345                    label: Some("deform test render"),
1346                    color_attachments: &[Some(wgpu::RenderPassColorAttachment {
1347                        view: &target.view,
1348                        resolve_target: None,
1349                        depth_slice: None,
1350                        ops: wgpu::Operations {
1351                            load: wgpu::LoadOp::Clear(wgpu::Color::GREEN),
1352                            store: wgpu::StoreOp::Store,
1353                        },
1354                    })],
1355                    depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
1356                        view: &depth_view,
1357                        depth_ops: Some(wgpu::Operations {
1358                            load: wgpu::LoadOp::Clear(1.0),
1359                            store: wgpu::StoreOp::Store,
1360                        }),
1361                        stencil_ops: None,
1362                    }),
1363                    timestamp_writes: None,
1364                    occlusion_query_set: None,
1365                });
1366                pass.set_pipeline(&pipeline);
1367                pass.set_bind_group(0, &empty_group, &[]);
1368                pass.set_bind_group(1, &empty_group, &[]);
1369                pass.set_bind_group(2, &deformers.group, &[]);
1370                pass.set_vertex_buffer(0, vertex_buffer.slice(..));
1371                pass.set_index_buffer(index_buffer.slice(..), wgpu::IndexFormat::Uint32);
1372                pass.draw_indexed(0..mesh.indices.len() as u32, 0, 0..1);
1373            }
1374            gpu.queue.submit(Some(encoder.finish()));
1375            let pixels = gpu.readback_rgba8(&target).unwrap();
1376            let mut raised_paper = 0;
1377            let mut raised_ink = 0;
1378            let mut flat_ink = 0;
1379            let mut flat_paper = 0;
1380            let mut shadow_mismatch = 0;
1381            for pixel in pixels.chunks_exact(4) {
1382                if pixel[0] > 200 && pixel[1] < 50 && pixel[2] > 200 {
1383                    shadow_mismatch += 1;
1384                }
1385                if pixel[0] < 50 && pixel[1] < 50 && pixel[2] < 50 {
1386                    if pixel[3] > 30 {
1387                        raised_ink += 1;
1388                    }
1389                    if pixel[3] == 0 {
1390                        flat_ink += 1;
1391                    }
1392                }
1393                if pixel[0] > 150 && pixel[1] > 150 && pixel[2] > 100 && pixel[3] > 30 {
1394                    raised_paper += 1;
1395                }
1396                if pixel[0] > 150 && pixel[1] > 150 && pixel[2] > 100 && pixel[3] == 0 {
1397                    flat_paper += 1;
1398                }
1399            }
1400            assert!(raised_paper > 100, "no raised paper at {at}");
1401            assert!(raised_ink > 10, "no attached ink at {at}");
1402            assert!(flat_paper > 100, "no flat paper at {at}");
1403            assert_eq!(shadow_mismatch, 0, "shadow deformation differs at {at}");
1404            if at > 0.5 {
1405                assert!(flat_ink > 10, "no ink on flat paper at {at}");
1406            }
1407        }
1408    }
1409}