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}