1use crate::chain::{Aux, Buffers, Chain, Pass, Uniforms};
2use crate::hash::BLUE_NOISE;
3use wgpu::{self, util::DeviceExt};
4
5const FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Rgba16Float;
6
7struct Target {
8 _texture: wgpu::Texture,
9 view: wgpu::TextureView,
10}
11
12impl Target {
13 fn new(device: &wgpu::Device, width: u32, height: u32) -> Self {
14 let texture = device.create_texture(&wgpu::TextureDescriptor {
15 label: Some("post intermediate"),
16 size: wgpu::Extent3d {
17 width,
18 height,
19 depth_or_array_layers: 1,
20 },
21 mip_level_count: 1,
22 sample_count: 1,
23 dimension: wgpu::TextureDimension::D2,
24 format: FORMAT,
25 usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::STORAGE_BINDING,
26 view_formats: &[],
27 });
28 let view = texture.create_view(&wgpu::TextureViewDescriptor::default());
29 Self {
30 _texture: texture,
31 view,
32 }
33 }
34}
35
36struct LowPair {
37 a: Target,
38 b: Target,
39}
40
41#[derive(Clone, Copy)]
42enum ImageSlot {
43 Main,
44 LowA,
45 LowB,
46}
47
48struct Stage {
49 pipeline: wgpu::ComputePipeline,
50 layout: wgpu::BindGroupLayout,
51 source: ImageSlot,
52 target: ImageSlot,
53 aux: Aux,
54 label: String,
55}
56
57struct Group {
58 pass: Pass,
59 stages: Vec<Stage>,
60 exposure: Vec<Option<usize>>,
61 low: Option<LowPair>,
62 factor: u32,
63 data: Option<wgpu::Buffer>,
64}
65
66pub struct GpuChain {
67 device: wgpu::Device,
68 groups: Vec<Group>,
69 tape: TapePlan,
70 scratch: [Target; 2],
71 bloom_sampler: Option<wgpu::Sampler>,
72 depth_buffer: wgpu::Buffer,
73 normal_fallback: Target,
74 depth_pipeline: wgpu::ComputePipeline,
75 width: u32,
76 height: u32,
77 viewport: (u32, u32),
78}
79
80struct Scissor<'a> {
81 rect: [u32; 4],
82 after: &'a [u32],
83}
84
85impl Scissor<'_> {
86 fn grown(&self, group: usize, chain: &GpuChain) -> [u32; 4] {
87 let reach = self.after[group];
88 [
89 self.rect[0].saturating_sub(reach),
90 self.rect[1].saturating_sub(reach),
91 (self.rect[2] + reach).min(chain.viewport.0),
92 (self.rect[3] + reach).min(chain.viewport.1),
93 ]
94 }
95}
96
97pub struct GpuStages {
98 pub bloom_grade: GpuChain,
99 pub final_stage: GpuChain,
100}
101
102impl GpuStages {
103 pub fn new(chain: &Chain, device: &wgpu::Device, width: u32, height: u32) -> Self {
104 let (bloom_grade, final_stage) = chain.split_after_bloom();
105 Self {
106 bloom_grade: GpuChain::new(bloom_grade, device, width, height),
107 final_stage: GpuChain::new(final_stage, device, width, height),
108 }
109 }
110
111 pub fn resize(&mut self, width: u32, height: u32) {
112 self.bloom_grade.resize(width, height);
113 self.final_stage.resize(width, height);
114 }
115}
116
117fn buffer(device: &wgpu::Device, label: &str, bytes: &[u8]) -> wgpu::Buffer {
118 device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
119 label: Some(label),
120 contents: bytes,
121 usage: wgpu::BufferUsages::STORAGE,
122 })
123}
124
125fn storage(device: &wgpu::Device, label: &str, size: u64) -> wgpu::Buffer {
126 device.create_buffer(&wgpu::BufferDescriptor {
127 label: Some(label),
128 size: size.max(4),
129 usage: wgpu::BufferUsages::STORAGE,
130 mapped_at_creation: false,
131 })
132}
133
134fn auxiliary_data(device: &wgpu::Device, pass: &Pass) -> Option<wgpu::Buffer> {
135 match pass {
136 Pass::Lut(lut) => {
137 let mut bytes = Vec::with_capacity(lut.values.len() * 12);
138 for rgb in &lut.values {
139 for value in rgb {
140 bytes.extend_from_slice(&value.to_ne_bytes());
141 }
142 }
143 if bytes.is_empty() {
144 bytes.extend_from_slice(&[0; 4]);
145 }
146 Some(buffer(device, "post lut", &bytes))
147 }
148 Pass::Grade(passes) => {
149 let bytes: Vec<u8> = crate::chain::grade_data(passes)
150 .iter()
151 .flat_map(|value| value.to_ne_bytes())
152 .collect();
153 Some(buffer(device, "post grade", &bytes))
154 }
155 Pass::OneBit(bit) if bit.kind == crate::fx::NoiseKind::Blue => {
156 let mut bytes = Vec::with_capacity(BLUE_NOISE.len() * 4);
157 for value in BLUE_NOISE {
158 bytes.extend_from_slice(&u32::from(*value).to_ne_bytes());
159 }
160 Some(buffer(device, "post blue noise", &bytes))
161 }
162 _ => None,
163 }
164}
165
166fn entry(binding: u32, ty: wgpu::BindingType) -> wgpu::BindGroupLayoutEntry {
167 wgpu::BindGroupLayoutEntry {
168 binding,
169 visibility: wgpu::ShaderStages::COMPUTE,
170 ty,
171 count: None,
172 }
173}
174
175fn sampled() -> wgpu::BindingType {
176 wgpu::BindingType::Texture {
177 sample_type: wgpu::TextureSampleType::Float { filterable: false },
178 view_dimension: wgpu::TextureViewDimension::D2,
179 multisampled: false,
180 }
181}
182
183fn storage_binding() -> wgpu::BindingType {
184 wgpu::BindingType::Buffer {
185 ty: wgpu::BufferBindingType::Storage { read_only: true },
186 has_dynamic_offset: false,
187 min_binding_size: None,
188 }
189}
190
191fn filterable() -> wgpu::BindingType {
192 wgpu::BindingType::Texture {
193 sample_type: wgpu::TextureSampleType::Float { filterable: true },
194 view_dimension: wgpu::TextureViewDimension::D2,
195 multisampled: false,
196 }
197}
198
199fn make_stage(
200 device: &wgpu::Device,
201 shader: &str,
202 aux: Aux,
203 source: ImageSlot,
204 target: ImageSlot,
205 label: String,
206) -> Stage {
207 let mut entries = vec![
208 entry(
209 0,
210 wgpu::BindingType::Buffer {
211 ty: wgpu::BufferBindingType::Uniform,
212 has_dynamic_offset: false,
213 min_binding_size: None,
214 },
215 ),
216 entry(
217 1,
218 if matches!(aux, Aux::LinearSampler) {
219 wgpu::BindingType::Texture {
220 sample_type: wgpu::TextureSampleType::Float { filterable: true },
221 view_dimension: wgpu::TextureViewDimension::D2,
222 multisampled: false,
223 }
224 } else {
225 sampled()
226 },
227 ),
228 entry(
229 2,
230 wgpu::BindingType::StorageTexture {
231 access: wgpu::StorageTextureAccess::WriteOnly,
232 format: FORMAT,
233 view_dimension: wgpu::TextureViewDimension::D2,
234 },
235 ),
236 ];
237 match aux {
238 Aux::Depth | Aux::Lut | Aux::Blue => entries.push(entry(3, storage_binding())),
239 Aux::Normal => entries.push(entry(4, sampled())),
240 Aux::DepthNormal => {
241 entries.push(entry(3, storage_binding()));
242 entries.push(entry(4, sampled()));
243 }
244 Aux::Bloom | Aux::BloomB => entries.push(entry(3, sampled())),
245 Aux::BloomLinear | Aux::GradeBloomLinear => {
246 entries.push(entry(3, filterable()));
247 entries.push(entry(
248 4,
249 wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering),
250 ));
251 }
252 Aux::GradeBloom | Aux::GradeBloomB => entries.push(entry(3, sampled())),
253 Aux::LinearSampler => entries.push(entry(
254 3,
255 wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering),
256 )),
257 Aux::Grade | Aux::None => {}
258 }
259 if matches!(
260 aux,
261 Aux::Grade | Aux::GradeBloom | Aux::GradeBloomB | Aux::GradeBloomLinear
262 ) {
263 entries.push(entry(5, storage_binding()));
264 }
265 let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
266 label: Some("post bindings"),
267 entries: &entries,
268 });
269 let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
270 label: Some("post pipeline layout"),
271 bind_group_layouts: &[Some(&layout)],
272 immediate_size: 0,
273 });
274 let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
275 label: Some(&label),
276 source: wgpu::ShaderSource::Wgsl(shader.into()),
277 });
278 let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
279 label: Some(&label),
280 layout: Some(&pipeline_layout),
281 module: &module,
282 entry_point: Some("main"),
283 compilation_options: wgpu::PipelineCompilationOptions::default(),
284 cache: None,
285 });
286 Stage {
287 pipeline,
288 layout,
289 source,
290 target,
291 aux,
292 label,
293 }
294}
295
296fn conversion_pipeline(device: &wgpu::Device) -> wgpu::ComputePipeline {
297 let texture = wgpu::BindingType::Texture {
298 sample_type: wgpu::TextureSampleType::Depth,
299 view_dimension: wgpu::TextureViewDimension::D2,
300 multisampled: false,
301 };
302 let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
303 label: Some("post geometry conversion"),
304 entries: &[
305 entry(0, texture),
306 entry(
307 1,
308 wgpu::BindingType::Buffer {
309 ty: wgpu::BufferBindingType::Storage { read_only: false },
310 has_dynamic_offset: false,
311 min_binding_size: None,
312 },
313 ),
314 entry(
315 2,
316 wgpu::BindingType::Buffer {
317 ty: wgpu::BufferBindingType::Uniform,
318 has_dynamic_offset: false,
319 min_binding_size: None,
320 },
321 ),
322 ],
323 });
324 let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
325 label: Some("post geometry conversion layout"),
326 bind_group_layouts: &[Some(&layout)],
327 immediate_size: 0,
328 });
329 let source = "@group(0) @binding(0) var src: texture_depth_2d; @group(0) @binding(1) var<storage, read_write> dst: array<f32>; @group(0) @binding(2) var<uniform> extent: vec4<u32>; @compute @workgroup_size(8, 8) fn main(@builtin(global_invocation_id) id: vec3<u32>) { let dim = min(textureDimensions(src), extent.xy); if (id.x < dim.x && id.y < dim.y) { dst[id.y * dim.x + id.x] = textureLoad(src, id.xy, 0); } }";
330 let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
331 label: Some("post geometry conversion shader"),
332 source: wgpu::ShaderSource::Wgsl(source.into()),
333 });
334 device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
335 label: Some("post geometry conversion pipeline"),
336 layout: Some(&pipeline_layout),
337 module: &module,
338 entry_point: Some("main"),
339 compilation_options: wgpu::PipelineCompilationOptions::default(),
340 cache: None,
341 })
342}
343
344fn plan_passes(passes: Vec<Pass>, fuse: bool) -> Vec<Pass> {
345 let active: Vec<Pass> = passes
346 .into_iter()
347 .filter(|pass| !pass.is_identity() || matches!(pass, Pass::Exposure(_)))
348 .collect();
349 let active = fuse_cavity_rim(active);
350 if fuse { fuse_grades(active) } else { active }
351}
352
353fn fuse_cavity_rim(passes: Vec<Pass>) -> Vec<Pass> {
354 let mut out: Vec<Pass> = Vec::with_capacity(passes.len());
355 for pass in passes {
356 if matches!(pass, Pass::Rim(_))
357 && matches!(out.last(), Some(Pass::Cavity(_)))
358 && let Some(cavity) = out.pop()
359 {
360 out.push(Pass::FusedCavityRim(vec![cavity, pass]));
361 } else {
362 out.push(pass);
363 }
364 }
365 out
366}
367
368fn joins_grade(run: &[Pass], pass: &Pass) -> bool {
369 match pass {
370 Pass::Bloom(_) => run.iter().all(|member| matches!(member, Pass::Exposure(_))),
371 other => other.is_per_pixel(),
372 }
373}
374
375fn fuse_grades(passes: Vec<Pass>) -> Vec<Pass> {
376 let mut out = Vec::with_capacity(passes.len());
377 let mut run: Vec<Pass> = Vec::new();
378 for pass in passes {
379 if !joins_grade(&run, &pass) {
380 close_grade(&mut out, std::mem::take(&mut run));
381 }
382 if joins_grade(&run, &pass) {
383 run.push(pass);
384 } else {
385 out.push(pass);
386 }
387 }
388 close_grade(&mut out, run);
389 out
390}
391
392fn close_grade(out: &mut Vec<Pass>, mut run: Vec<Pass>) {
393 match run.len() {
394 0 => {}
395 1 => out.extend(run.pop()),
396 _ => out.push(Pass::Grade(run)),
397 }
398}
399
400fn exposure_slots(pass: &Pass, stages: usize) -> Vec<Option<usize>> {
401 let mut slots = vec![None; stages];
402 match pass {
403 Pass::Exposure(_) => slots[0] = Some(0),
404 Pass::Grade(passes) if matches!(passes.first(), Some(Pass::Exposure(_))) => {
405 if let Some(last) = slots.last_mut() {
406 *last = Some(0);
407 }
408 match passes
409 .iter()
410 .find(|member| !matches!(member, Pass::Exposure(_)))
411 {
412 Some(Pass::Bloom(bloom)) if bloom.kind == crate::bloom::BloomKind::Rings => {
413 slots[0] = Some(3)
414 }
415 Some(Pass::Bloom(_)) => slots[0] = Some(6),
416 _ => {}
417 }
418 }
419 _ => {}
420 }
421 slots
422}
423
424fn stage_name(pass: &Pass, stage: usize, count: usize) -> String {
425 let parts: &[&str] = match count {
426 3 => &["extract", "rings", "add"],
427 4 => &["extract", "blur x", "blur y", "add"],
428 _ => &[],
429 };
430 match parts.get(stage) {
431 Some(part) => format!("{} {part}", pass.name()),
432 None => pass.name().to_owned(),
433 }
434}
435
436fn plan_groups(
437 passes: Vec<Pass>,
438 fuse: bool,
439 device: &wgpu::Device,
440 width: u32,
441 height: u32,
442) -> Vec<Group> {
443 let active = plan_passes(passes, fuse);
444 let passes = if active.is_empty() {
445 vec![Pass::Exposure(1.0)]
446 } else {
447 active
448 };
449 let mut groups = Vec::with_capacity(passes.len());
450 for (index, pass) in passes.into_iter().enumerate() {
451 let dispatches = pass.dispatches(Buffers::GEOMETRY);
452 let bloom = match &pass {
453 Pass::Grade(passes) => passes.iter().find_map(|member| match member {
454 Pass::Bloom(bloom) => Some(*bloom),
455 _ => None,
456 }),
457 _ => None,
458 };
459 let (factor, slots): (u32, Vec<(ImageSlot, ImageSlot)>) = match (&pass, bloom) {
460 (Pass::Grade(_), Some(bloom)) if dispatches.len() == 3 => (
461 bloom.down.max(1),
462 vec![
463 (ImageSlot::Main, ImageSlot::LowA),
464 (ImageSlot::LowA, ImageSlot::LowB),
465 (ImageSlot::Main, ImageSlot::Main),
466 ],
467 ),
468 (Pass::Grade(_), Some(bloom)) => (
469 bloom.down.max(1),
470 vec![
471 (ImageSlot::Main, ImageSlot::LowA),
472 (ImageSlot::LowA, ImageSlot::LowB),
473 (ImageSlot::LowB, ImageSlot::LowA),
474 (ImageSlot::Main, ImageSlot::Main),
475 ],
476 ),
477 (Pass::Bloom(bloom), _) if bloom.kind == crate::bloom::BloomKind::Rings => (
478 bloom.down.max(1),
479 vec![
480 (ImageSlot::Main, ImageSlot::LowA),
481 (ImageSlot::LowA, ImageSlot::LowB),
482 (ImageSlot::Main, ImageSlot::Main),
483 ],
484 ),
485 (Pass::Bloom(bloom), _) if dispatches.len() == 4 => (
486 bloom.down.max(1),
487 vec![
488 (ImageSlot::Main, ImageSlot::LowA),
489 (ImageSlot::LowA, ImageSlot::LowB),
490 (ImageSlot::LowB, ImageSlot::LowA),
491 (ImageSlot::Main, ImageSlot::Main),
492 ],
493 ),
494 (Pass::Halation(_), _) => (
495 4,
496 vec![
497 (ImageSlot::Main, ImageSlot::LowA),
498 (ImageSlot::LowA, ImageSlot::LowB),
499 (ImageSlot::LowB, ImageSlot::LowA),
500 (ImageSlot::Main, ImageSlot::Main),
501 ],
502 ),
503 (Pass::Neon(_), _) if dispatches.len() == 4 => (
504 2,
505 vec![
506 (ImageSlot::Main, ImageSlot::LowA),
507 (ImageSlot::LowA, ImageSlot::LowB),
508 (ImageSlot::LowB, ImageSlot::LowA),
509 (ImageSlot::Main, ImageSlot::Main),
510 ],
511 ),
512 _ => (1, vec![(ImageSlot::Main, ImageSlot::Main)]),
513 };
514 let stages = dispatches
515 .iter()
516 .zip(slots)
517 .enumerate()
518 .map(|(stage_index, (dispatch, (source, target)))| {
519 let label = format!(
520 "post {index}:{stage_index} {}",
521 stage_name(&pass, stage_index, dispatches.len())
522 );
523 make_stage(
524 device,
525 &dispatch.shader,
526 dispatch.aux,
527 source,
528 target,
529 label,
530 )
531 })
532 .collect();
533 let low = if (bloom.is_some()
534 || matches!(pass, Pass::Bloom(_) | Pass::Halation(_) | Pass::Neon(_)))
535 && dispatches.len() > 1
536 {
537 let w = (width / factor).max(1);
538 let h = (height / factor).max(1);
539 Some(LowPair {
540 a: Target::new(device, w, h),
541 b: Target::new(device, w, h),
542 })
543 } else {
544 None
545 };
546 let data = auxiliary_data(device, &pass);
547 let exposure = exposure_slots(&pass, dispatches.len());
548 groups.push(Group {
549 pass,
550 stages,
551 exposure,
552 low,
553 factor,
554 data,
555 });
556 }
557 groups
558}
559
560struct TapePlan {
561 source: Vec<Pass>,
562 fuse: bool,
563 groups: Vec<Group>,
564 tape: crate::tape::Tape,
565 encodes: bool,
566}
567
568impl TapePlan {
569 fn plan(&mut self, device: &wgpu::Device, width: u32, height: u32) {
570 let mut chain = Chain {
571 passes: self.source.clone(),
572 seed: 0,
573 frame: 0,
574 };
575 chain.set_tape(self.tape);
576 let mut groups = plan_groups(chain.passes, self.fuse, device, width, height);
577 self.encodes = groups.len() >= 2
578 && matches!(groups[groups.len() - 2].pass, Pass::Tape(_))
579 && matches!(groups[groups.len() - 1].pass, Pass::Encode);
580 if self.encodes {
581 groups.pop();
582 }
583 self.groups = groups;
584 }
585
586 fn set(&mut self, tape: crate::tape::Tape) {
587 self.tape = tape;
588 for group in &mut self.groups {
589 if matches!(group.pass, Pass::Tape(_)) {
590 group.pass = Pass::Tape(tape);
591 }
592 }
593 }
594}
595
596pub(crate) fn fit_viewport(
597 viewport: (u32, u32),
598 capacity: (u32, u32),
599) -> Result<(u32, u32), String> {
600 let (width, height) = viewport;
601 if width == 0 || height == 0 {
602 return Err(format!("a viewport of {width}x{height} is empty"));
603 }
604 if width > capacity.0 || height > capacity.1 {
605 return Err(format!(
606 "a viewport of {width}x{height} does not fit the {}x{} allocation",
607 capacity.0, capacity.1
608 ));
609 }
610 Ok(viewport)
611}
612
613impl GpuChain {
614 pub fn new(chain: Chain, device: &wgpu::Device, width: u32, height: u32) -> Self {
615 Self::build(chain, device, width, height, true)
616 }
617
618 #[cfg(test)]
619 fn unfused(chain: Chain, device: &wgpu::Device, width: u32, height: u32) -> Self {
620 Self::build(chain, device, width, height, false)
621 }
622
623 fn build(chain: Chain, device: &wgpu::Device, width: u32, height: u32, fuse: bool) -> Self {
624 assert!(width > 0 && height > 0);
625 let taped = chain.tape();
626 let plain = chain
627 .passes
628 .iter()
629 .filter(|pass| !matches!(pass, Pass::Tape(_)))
630 .cloned()
631 .collect();
632 let groups = plan_groups(plain, fuse, device, width, height);
633 let mut tape = TapePlan {
634 source: chain.passes,
635 fuse,
636 groups: Vec::new(),
637 tape: crate::tape::Tape::OFF,
638 encodes: false,
639 };
640 if let Some(taped) = taped {
641 tape.tape = taped;
642 tape.plan(device, width, height);
643 }
644 let bloom_sampler = groups
645 .iter()
646 .chain(&tape.groups)
647 .flat_map(|group| &group.stages)
648 .any(|stage| {
649 matches!(
650 stage.aux,
651 Aux::BloomLinear | Aux::LinearSampler | Aux::GradeBloomLinear
652 )
653 })
654 .then(|| {
655 device.create_sampler(&wgpu::SamplerDescriptor {
656 mag_filter: wgpu::FilterMode::Linear,
657 min_filter: wgpu::FilterMode::Linear,
658 ..Default::default()
659 })
660 });
661 let pixels = u64::from(width) * u64::from(height);
662 Self {
663 device: device.clone(),
664 groups,
665 tape,
666 scratch: [
667 Target::new(device, width, height),
668 Target::new(device, width, height),
669 ],
670 bloom_sampler,
671 depth_buffer: storage(device, "post depth", pixels * 4),
672 normal_fallback: Target::new(device, 1, 1),
673 depth_pipeline: conversion_pipeline(device),
674 width,
675 height,
676 viewport: (width, height),
677 }
678 }
679
680 pub fn size(&self) -> (u32, u32) {
681 (self.width, self.height)
682 }
683
684 pub fn capacity(&self) -> (u32, u32) {
685 (self.width, self.height)
686 }
687
688 pub fn viewport(&self) -> (u32, u32) {
689 self.viewport
690 }
691
692 pub fn set_viewport(&mut self, width: u32, height: u32) -> Result<(), String> {
693 self.viewport = fit_viewport((width, height), (self.width, self.height))?;
694 Ok(())
695 }
696
697 fn low_size(&self, group: &Group) -> (u32, u32) {
698 (
699 (self.viewport.0 / group.factor).max(1),
700 (self.viewport.1 / group.factor).max(1),
701 )
702 }
703
704 fn slot_size(&self, group: &Group, slot: ImageSlot) -> (u32, u32) {
705 match slot {
706 ImageSlot::Main => self.viewport,
707 ImageSlot::LowA | ImageSlot::LowB => self.low_size(group),
708 }
709 }
710
711 pub fn tape(&self) -> crate::tape::Tape {
712 self.tape.tape
713 }
714
715 pub fn set_tape(&mut self, tape: crate::tape::Tape) {
716 if self.tape.groups.is_empty() && !tape.is_off() {
717 self.tape.tape = tape;
718 self.tape.plan(&self.device, self.width, self.height);
719 self.bloom_sampler.get_or_insert_with(|| {
720 self.device.create_sampler(&wgpu::SamplerDescriptor {
721 mag_filter: wgpu::FilterMode::Linear,
722 min_filter: wgpu::FilterMode::Linear,
723 ..Default::default()
724 })
725 });
726 }
727 self.tape.set(tape);
728 }
729
730 pub fn reset_tape(&mut self) {
731 let own = self
732 .tape
733 .source
734 .iter()
735 .find_map(|pass| match pass {
736 Pass::Tape(tape) => Some(*tape),
737 _ => None,
738 })
739 .unwrap_or(crate::tape::Tape::OFF);
740 self.set_tape(own);
741 }
742
743 fn active(&self) -> &[Group] {
744 if self.tape.tape.is_off() || self.tape.groups.is_empty() {
745 &self.groups
746 } else {
747 &self.tape.groups
748 }
749 }
750
751 pub fn resize(&mut self, width: u32, height: u32) {
752 if width == 0 || height == 0 {
753 return;
754 }
755 self.viewport = (width, height);
756 if (width, height) == (self.width, self.height) {
757 return;
758 }
759 self.width = width;
760 self.height = height;
761 self.scratch = [
762 Target::new(&self.device, width, height),
763 Target::new(&self.device, width, height),
764 ];
765 let pixels = u64::from(width) * u64::from(height);
766 self.depth_buffer = storage(&self.device, "post depth", pixels * 4);
767 for group in self.groups.iter_mut().chain(&mut self.tape.groups) {
768 if group.low.is_some() {
769 let w = (width / group.factor).max(1);
770 let h = (height / group.factor).max(1);
771 group.low = Some(LowPair {
772 a: Target::new(&self.device, w, h),
773 b: Target::new(&self.device, w, h),
774 });
775 }
776 }
777 }
778
779 fn convert(
780 &self,
781 encoder: &mut wgpu::CommandEncoder,
782 view: &wgpu::TextureView,
783 mut profiler: Option<&mut pfx_gpu::GpuProfiler>,
784 ) {
785 let pipeline = &self.depth_pipeline;
786 let dst = &self.depth_buffer;
787 let (width, height) = self.viewport;
788 let extent = self
789 .device
790 .create_buffer_init(&wgpu::util::BufferInitDescriptor {
791 label: Some("post geometry conversion extent"),
792 contents: bytemuck::cast_slice(&[width, height, 0, 0]),
793 usage: wgpu::BufferUsages::UNIFORM,
794 });
795 let layout = pipeline.get_bind_group_layout(0);
796 let bind = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
797 label: Some("post geometry conversion bindings"),
798 layout: &layout,
799 entries: &[
800 wgpu::BindGroupEntry {
801 binding: 0,
802 resource: wgpu::BindingResource::TextureView(view),
803 },
804 wgpu::BindGroupEntry {
805 binding: 1,
806 resource: dst.as_entire_binding(),
807 },
808 wgpu::BindGroupEntry {
809 binding: 2,
810 resource: extent.as_entire_binding(),
811 },
812 ],
813 });
814 let timing = profiler
815 .as_deref_mut()
816 .and_then(|p| p.pass("post depth conversion"));
817 let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
818 label: Some("post geometry conversion"),
819 timestamp_writes: profiler.as_ref().and_then(|p| p.compute_writes(timing)),
820 });
821 pass.set_pipeline(pipeline);
822 pass.set_bind_group(0, &bind, &[]);
823 pass.dispatch_workgroups(width.div_ceil(8), height.div_ceil(8), 1);
824 }
825
826 #[allow(clippy::too_many_arguments)]
827 pub fn run(
828 &self,
829 encoder: &mut wgpu::CommandEncoder,
830 input: &wgpu::TextureView,
831 depth: Option<&wgpu::TextureView>,
832 normal: Option<&wgpu::TextureView>,
833 output: &wgpu::TextureView,
834 frame: u32,
835 seed: u32,
836 ) {
837 self.run_timed(encoder, input, depth, normal, output, frame, seed, None);
838 }
839
840 #[allow(clippy::too_many_arguments)]
841 pub fn run_timed(
842 &self,
843 encoder: &mut wgpu::CommandEncoder,
844 input: &wgpu::TextureView,
845 depth: Option<&wgpu::TextureView>,
846 normal: Option<&wgpu::TextureView>,
847 output: &wgpu::TextureView,
848 frame: u32,
849 seed: u32,
850 mut profiler: Option<&mut pfx_gpu::GpuProfiler>,
851 ) {
852 self.run_timed_exposed(
853 encoder,
854 input,
855 depth,
856 normal,
857 output,
858 frame,
859 seed,
860 1.0,
861 profiler.take(),
862 );
863 }
864
865 #[allow(clippy::too_many_arguments)]
866 pub fn run_timed_exposed(
867 &self,
868 encoder: &mut wgpu::CommandEncoder,
869 input: &wgpu::TextureView,
870 depth: Option<&wgpu::TextureView>,
871 normal: Option<&wgpu::TextureView>,
872 output: &wgpu::TextureView,
873 frame: u32,
874 seed: u32,
875 exposure: f32,
876 profiler: Option<&mut pfx_gpu::GpuProfiler>,
877 ) {
878 self.encode(
879 encoder, input, depth, normal, output, frame, seed, exposure, None, profiler,
880 );
881 }
882
883 pub fn region_reach(&self) -> Option<u32> {
884 let groups = self.active();
885 let mut total = 0u32;
886 for group in groups {
887 let local = group.stages.len() == 1
888 && matches!(group.stages[0].source, ImageSlot::Main)
889 && matches!(group.stages[0].target, ImageSlot::Main)
890 && !matches!(
891 group.stages[0].aux,
892 Aux::Depth | Aux::DepthNormal | Aux::LinearSampler
893 );
894 if !local {
895 return None;
896 }
897 total += group.pass.region_reach()?;
898 }
899 Some(total)
900 }
901
902 #[allow(clippy::too_many_arguments)]
903 pub fn run_region(
904 &self,
905 encoder: &mut wgpu::CommandEncoder,
906 input: &wgpu::TextureView,
907 depth: Option<&wgpu::TextureView>,
908 normal: Option<&wgpu::TextureView>,
909 output: &wgpu::TextureView,
910 frame: u32,
911 seed: u32,
912 exposure: f32,
913 rects: &[[u32; 4]],
914 mut profiler: Option<&mut pfx_gpu::GpuProfiler>,
915 ) -> bool {
916 let rects: Vec<[u32; 4]> = rects
917 .iter()
918 .map(|&[x0, y0, x1, y1]| [x0, y0, x1.min(self.viewport.0), y1.min(self.viewport.1)])
919 .filter(|rect| rect[0] < rect[2] && rect[1] < rect[3])
920 .collect();
921 let Some(total) = self.region_reach() else {
922 self.encode(
923 encoder, input, depth, normal, output, frame, seed, exposure, None, profiler,
924 );
925 return false;
926 };
927 let mut after = Vec::with_capacity(self.active().len());
928 let mut left = total;
929 for group in self.active() {
930 left -= group.pass.region_reach().unwrap_or(0);
931 after.push(left);
932 }
933 for rect in rects {
934 let scissor = Scissor {
935 rect,
936 after: &after,
937 };
938 self.encode(
939 encoder,
940 input,
941 depth,
942 normal,
943 output,
944 frame,
945 seed,
946 exposure,
947 Some(scissor),
948 profiler.as_deref_mut(),
949 );
950 }
951 true
952 }
953
954 #[allow(clippy::too_many_arguments)]
955 fn encode(
956 &self,
957 encoder: &mut wgpu::CommandEncoder,
958 input: &wgpu::TextureView,
959 depth: Option<&wgpu::TextureView>,
960 normal: Option<&wgpu::TextureView>,
961 output: &wgpu::TextureView,
962 frame: u32,
963 seed: u32,
964 exposure: f32,
965 scissor: Option<Scissor<'_>>,
966 mut profiler: Option<&mut pfx_gpu::GpuProfiler>,
967 ) {
968 let groups = self.active();
969 if let Some(view) = depth
970 && groups.iter().any(|group| {
971 group
972 .stages
973 .iter()
974 .any(|stage| matches!(stage.aux, Aux::Depth | Aux::DepthNormal))
975 })
976 {
977 self.convert(encoder, view, profiler.as_deref_mut());
978 }
979 let buffers = Buffers {
980 depth: depth.is_some(),
981 normal: normal.is_some(),
982 };
983 let mut current = input;
984 let mut scratch_index = 0;
985 for (group_index, group) in groups.iter().enumerate() {
986 let mut params = group.pass.params(buffers);
987 if self.tape.encodes && matches!(group.pass, Pass::Tape(_)) {
988 params[0][14] = 1.0;
989 }
990 for (stage_index, (stage, mut params)) in group.stages.iter().zip(params).enumerate() {
991 let last_stage = stage_index + 1 == group.stages.len();
992 let last_group = group_index + 1 == groups.len();
993 let target = match stage.target {
994 ImageSlot::Main if last_stage && last_group => output,
995 ImageSlot::Main => {
996 let view = &self.scratch[scratch_index].view;
997 scratch_index ^= 1;
998 view
999 }
1000 ImageSlot::LowA => &group.low.as_ref().unwrap().a.view,
1001 ImageSlot::LowB => &group.low.as_ref().unwrap().b.view,
1002 };
1003 let source = match stage.source {
1004 ImageSlot::Main => current,
1005 ImageSlot::LowA => &group.low.as_ref().unwrap().a.view,
1006 ImageSlot::LowB => &group.low.as_ref().unwrap().b.view,
1007 };
1008 let (w, h) = self.slot_size(group, stage.target);
1009 if group_index == 0
1010 && let Some(slot) = group.exposure[stage_index]
1011 {
1012 params[slot] *= exposure;
1013 }
1014 let area = match (&scissor, stage.target) {
1015 (Some(scissor), ImageSlot::Main) => Some(scissor.grown(group_index, self)),
1016 _ => None,
1017 };
1018 let mut uniforms = Uniforms::new(w, h, frame, seed, params)
1019 .extent(self.slot_size(group, stage.source), self.low_size(group));
1020 if let Some(area) = area {
1021 uniforms = uniforms.region(area);
1022 }
1023 let uniform = self
1024 .device
1025 .create_buffer_init(&wgpu::util::BufferInitDescriptor {
1026 label: Some("post uniforms"),
1027 contents: &uniforms.bytes(),
1028 usage: wgpu::BufferUsages::UNIFORM,
1029 });
1030 let mut entries = vec![
1031 wgpu::BindGroupEntry {
1032 binding: 0,
1033 resource: uniform.as_entire_binding(),
1034 },
1035 wgpu::BindGroupEntry {
1036 binding: 1,
1037 resource: wgpu::BindingResource::TextureView(source),
1038 },
1039 wgpu::BindGroupEntry {
1040 binding: 2,
1041 resource: wgpu::BindingResource::TextureView(target),
1042 },
1043 ];
1044 match stage.aux {
1045 Aux::Depth => entries.push(wgpu::BindGroupEntry {
1046 binding: 3,
1047 resource: self.depth_buffer.as_entire_binding(),
1048 }),
1049 Aux::Normal => entries.push(wgpu::BindGroupEntry {
1050 binding: 4,
1051 resource: wgpu::BindingResource::TextureView(
1052 normal.unwrap_or(&self.normal_fallback.view),
1053 ),
1054 }),
1055 Aux::DepthNormal => {
1056 entries.push(wgpu::BindGroupEntry {
1057 binding: 3,
1058 resource: self.depth_buffer.as_entire_binding(),
1059 });
1060 entries.push(wgpu::BindGroupEntry {
1061 binding: 4,
1062 resource: wgpu::BindingResource::TextureView(
1063 normal.unwrap_or(&self.normal_fallback.view),
1064 ),
1065 });
1066 }
1067 Aux::Bloom | Aux::BloomB => entries.push(wgpu::BindGroupEntry {
1068 binding: 3,
1069 resource: wgpu::BindingResource::TextureView(
1070 if matches!(stage.aux, Aux::BloomB) {
1071 &group.low.as_ref().unwrap().b.view
1072 } else {
1073 &group.low.as_ref().unwrap().a.view
1074 },
1075 ),
1076 }),
1077 Aux::BloomLinear => {
1078 entries.push(wgpu::BindGroupEntry {
1079 binding: 3,
1080 resource: wgpu::BindingResource::TextureView(
1081 &group.low.as_ref().unwrap().a.view,
1082 ),
1083 });
1084 entries.push(wgpu::BindGroupEntry {
1085 binding: 4,
1086 resource: wgpu::BindingResource::Sampler(
1087 self.bloom_sampler.as_ref().unwrap(),
1088 ),
1089 });
1090 }
1091 Aux::LinearSampler => entries.push(wgpu::BindGroupEntry {
1092 binding: 3,
1093 resource: wgpu::BindingResource::Sampler(
1094 self.bloom_sampler.as_ref().unwrap(),
1095 ),
1096 }),
1097 Aux::Lut | Aux::Blue => entries.push(wgpu::BindGroupEntry {
1098 binding: 3,
1099 resource: group.data.as_ref().unwrap().as_entire_binding(),
1100 }),
1101 Aux::GradeBloom | Aux::GradeBloomLinear => entries.push(wgpu::BindGroupEntry {
1102 binding: 3,
1103 resource: wgpu::BindingResource::TextureView(
1104 &group.low.as_ref().unwrap().a.view,
1105 ),
1106 }),
1107 Aux::GradeBloomB => entries.push(wgpu::BindGroupEntry {
1108 binding: 3,
1109 resource: wgpu::BindingResource::TextureView(
1110 &group.low.as_ref().unwrap().b.view,
1111 ),
1112 }),
1113 Aux::Grade | Aux::None => {}
1114 }
1115 if stage.aux == Aux::GradeBloomLinear {
1116 entries.push(wgpu::BindGroupEntry {
1117 binding: 4,
1118 resource: wgpu::BindingResource::Sampler(
1119 self.bloom_sampler.as_ref().unwrap(),
1120 ),
1121 });
1122 }
1123 if matches!(
1124 stage.aux,
1125 Aux::Grade | Aux::GradeBloom | Aux::GradeBloomB | Aux::GradeBloomLinear
1126 ) {
1127 entries.push(wgpu::BindGroupEntry {
1128 binding: 5,
1129 resource: group.data.as_ref().unwrap().as_entire_binding(),
1130 });
1131 }
1132 let bind = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
1133 label: Some(&stage.label),
1134 layout: &stage.layout,
1135 entries: &entries,
1136 });
1137 let timing = profiler
1138 .as_deref_mut()
1139 .and_then(|timer| timer.pass(stage.label.clone()));
1140 let writes = profiler
1141 .as_deref()
1142 .and_then(|timer| timer.compute_writes(timing));
1143 let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
1144 label: Some(&stage.label),
1145 timestamp_writes: writes,
1146 });
1147 pass.set_pipeline(&stage.pipeline);
1148 pass.set_bind_group(0, &bind, &[]);
1149 let (groups_x, groups_y) = match area {
1150 Some([x0, y0, x1, y1]) => ((x1 - x0).div_ceil(8), (y1 - y0).div_ceil(8)),
1151 None => (w.div_ceil(8), h.div_ceil(8)),
1152 };
1153 pass.dispatch_workgroups(groups_x, groups_y, 1);
1154 drop(pass);
1155 if matches!(stage.target, ImageSlot::Main) {
1156 current = target;
1157 }
1158 }
1159 }
1160 }
1161}
1162
1163#[cfg(test)]
1164pub(crate) mod tests {
1165 use super::*;
1166 use crate::image::Image;
1167 use crate::preset::Style;
1168 use pfx_gpu::{Gpu, OffscreenTarget};
1169
1170 pub(crate) fn half(value: f32) -> u16 {
1171 let bits = value.to_bits();
1172 let sign = ((bits >> 16) & 0x8000) as u16;
1173 let exponent = ((bits >> 23) & 255) as i32 - 127 + 15;
1174 let mantissa = bits & 0x7fffff;
1175 if exponent <= 0 {
1176 return sign;
1177 }
1178 if exponent >= 31 {
1179 return sign | 0x7c00;
1180 }
1181 sign | ((exponent as u16) << 10) | ((mantissa >> 13) as u16)
1182 }
1183
1184 fn float(value: u16) -> f32 {
1185 let sign = if value & 0x8000 != 0 { -1.0 } else { 1.0 };
1186 let exponent = ((value >> 10) & 31) as i32;
1187 let mantissa = f32::from(value & 1023) / 1024.0;
1188 if exponent == 0 {
1189 sign * mantissa * 2f32.powi(-14)
1190 } else if exponent == 31 {
1191 f32::INFINITY
1192 } else {
1193 sign * (1.0 + mantissa) * 2f32.powi(exponent - 15)
1194 }
1195 }
1196
1197 fn synthetic(width: u32, height: u32) -> Image {
1198 Image::from_fn(width, height, |x, y| {
1199 let shades = [0.0, 0.25, 0.5, 0.75, 1.0, 2.0];
1200 [
1201 shades[((x / 8 + y / 16) % 6) as usize],
1202 shades[((x / 16 + y / 8 + 2) % 6) as usize],
1203 shades[((x / 8 + 3) % 6) as usize],
1204 1.0,
1205 ]
1206 })
1207 }
1208
1209 fn output_target(gpu: &Gpu, width: u32, height: u32) -> OffscreenTarget {
1210 let texture = gpu.device.create_texture(&wgpu::TextureDescriptor {
1211 label: Some("post test output"),
1212 size: wgpu::Extent3d {
1213 width,
1214 height,
1215 depth_or_array_layers: 1,
1216 },
1217 mip_level_count: 1,
1218 sample_count: 1,
1219 dimension: wgpu::TextureDimension::D2,
1220 format: FORMAT,
1221 usage: wgpu::TextureUsages::STORAGE_BINDING
1222 | wgpu::TextureUsages::TEXTURE_BINDING
1223 | wgpu::TextureUsages::COPY_SRC
1224 | wgpu::TextureUsages::COPY_DST,
1225 view_formats: &[],
1226 });
1227 let view = texture.create_view(&wgpu::TextureViewDescriptor::default());
1228 OffscreenTarget {
1229 texture,
1230 view,
1231 format: FORMAT,
1232 width,
1233 height,
1234 }
1235 }
1236
1237 fn plan_names(passes: Vec<Pass>) -> Vec<String> {
1238 plan_passes(passes, true)
1239 .iter()
1240 .map(|pass| match pass {
1241 Pass::Grade(members) => format!(
1242 "grade({})",
1243 members.iter().map(Pass::name).collect::<Vec<_>>().join("+")
1244 ),
1245 other => other.name().to_owned(),
1246 })
1247 .collect()
1248 }
1249
1250 fn fused_finish(warmth: f32, vignette: f32) -> Vec<Pass> {
1251 vec![
1252 Pass::Exposure(1.3),
1253 Pass::Bloom(crate::bloom::Bloom::tent(0.25, 0.5, 4, 16.0, 8)),
1254 Pass::Black(0.02),
1255 Pass::Tone(crate::tone::Tone::neutral(0.75, false)),
1256 Pass::Warmth(crate::grade::Warmth::luma_curve(
1257 warmth,
1258 [1.0, 1.0],
1259 [1.02, 0.98],
1260 0.2,
1261 0.8,
1262 )),
1263 Pass::Vignette(crate::grade::Vignette::smoothstep(vignette, 1.0, 0.2, 0.8)),
1264 Pass::Grain(crate::grade::Grain::clamped(0.02, 0.5)),
1265 ]
1266 }
1267
1268 #[test]
1269 fn fused_finish_is_one_grade() {
1270 for (warmth, vignette, names) in [
1271 (
1272 0.5,
1273 0.3,
1274 "grade(exposure+bloom+black+tone+warmth+vignette+grain)",
1275 ),
1276 (0.0, 0.3, "grade(exposure+bloom+black+tone+vignette+grain)"),
1277 (0.7, 0.0, "grade(exposure+bloom+black+tone+warmth+grain)"),
1278 (0.0, 0.0, "grade(exposure+bloom+black+tone+grain)"),
1279 ] {
1280 assert_eq!(
1281 plan_names(fused_finish(warmth, vignette)),
1282 [names],
1283 "warmth {warmth} vignette {vignette}"
1284 );
1285 }
1286 }
1287
1288 #[test]
1289 fn identity_passes_leave_the_grade() {
1290 let passes = vec![
1291 Pass::Exposure(1.0),
1292 Pass::Tone(crate::tone::Tone::neutral(1.0, false)),
1293 Pass::Warmth(crate::grade::Warmth::luma_curve(
1294 0.0, [1.0; 2], [1.0; 2], 0.2, 0.8,
1295 )),
1296 Pass::Grain(crate::grade::Grain::clamped(0.02, 0.5)),
1297 Pass::Warmth(crate::grade::Warmth::linear(0.0)),
1298 Pass::Vignette(crate::grade::Vignette::smoothstep(0.0, 1.0, 0.2, 0.8)),
1299 ];
1300 assert_eq!(plan_names(passes), ["grade(exposure+tone+grain)"]);
1301 }
1302
1303 #[test]
1304 fn a_grade_takes_one_bloom_after_exposure_only() {
1305 let bloom = Pass::Bloom(crate::bloom::Bloom::neutral(0.2));
1306 let passes = vec![
1307 Pass::Exposure(1.0),
1308 Pass::Exposure(0.9),
1309 bloom.clone(),
1310 Pass::Vignette(crate::grade::Vignette::power(0.4, 2.0, 1.0, 1.0)),
1311 bloom.clone(),
1312 Pass::Encode,
1313 ];
1314 assert_eq!(
1315 plan_names(passes),
1316 [
1317 "grade(exposure+exposure+bloom+vignette)",
1318 "grade(bloom+encode)"
1319 ]
1320 );
1321 let passes = vec![
1322 Pass::Tone(crate::tone::Tone::aces()),
1323 bloom.clone(),
1324 Pass::Kuwahara(2),
1325 Pass::Encode,
1326 ];
1327 assert_eq!(plan_names(passes), ["tone", "bloom", "kuwahara", "encode"]);
1328 let unfused = plan_passes(vec![Pass::Exposure(0.9), bloom, Pass::Encode], false);
1329 assert_eq!(
1330 unfused.iter().map(Pass::name).collect::<Vec<_>>(),
1331 ["exposure", "bloom", "encode"]
1332 );
1333 }
1334
1335 #[test]
1336 fn styles_fuse_their_per_pixel_runs() {
1337 let expected: [(&str, &[&str]); 21] = [
1338 ("cel", &["exposure", "cel", "fused_cavity_rim", "outline"]),
1339 ("bw", &["bw"]),
1340 ("noir", &["grade(noir+grain+vignette)"]),
1341 ("vignette", &["vignette"]),
1342 ("neon", &["neon", "bloom"]),
1343 ("rubber_hose", &["rubber_hose"]),
1344 ("sepia", &["sepia"]),
1345 (
1346 "film",
1347 &[
1348 "distortion",
1349 "aberration",
1350 "halation",
1351 "grade(tone+sepia+grain+vignette)",
1352 ],
1353 ),
1354 (
1355 "crt",
1356 &[
1357 "curvature",
1358 "aberration",
1359 "grade(bloom+scanlines+aperture+vignette)",
1360 ],
1361 ),
1362 ("one_bit", &["one_bit"]),
1363 ("halftone", &["halftone"]),
1364 ("comic", &["comic"]),
1365 ("modern_comic", &["modern_comic"]),
1366 ("watercolor", &["watercolor"]),
1367 ("paper_grain", &["paper_grain"]),
1368 ("pixel", &["pixel"]),
1369 ("duotone", &["duotone"]),
1370 ("gradient_map", &["gradient_map"]),
1371 ("posterize", &["posterize"]),
1372 ("kuwahara", &["kuwahara"]),
1373 ("tilt_shift", &["tilt_shift"]),
1374 ];
1375 for (style, (name, names)) in Style::ALL.into_iter().zip(expected) {
1376 assert_eq!(style.name(), name);
1377 assert_eq!(plan_names(style.chain().passes), names, "{name}");
1378 }
1379 assert_eq!(
1380 plan_names(as_live(frame_finish()).passes),
1381 ["grade(exposure+exposure+bloom+vignette+saturation+tone+encode)"]
1382 );
1383 }
1384
1385 #[test]
1386 fn frame_finish_is_four_dispatches_with_its_exposure_in_two() {
1387 let plan = plan_passes(as_live(frame_finish()).passes, true);
1388 let [Pass::Grade(members)] = plan.as_slice() else {
1389 panic!("{plan:?}");
1390 };
1391 let dispatches = plan[0].dispatches(Buffers::GEOMETRY);
1392 assert_eq!(dispatches.len(), 4);
1393 assert_eq!(dispatches[0].shader, crate::wgsl::BLOOM_DOWN_BOX2);
1394 assert_eq!(
1395 dispatches[1].shader.as_ref(),
1396 crate::wgsl::bloom_blur_linear(4, 16.0)
1397 );
1398 assert_eq!(dispatches[3].aux, Aux::GradeBloomLinear);
1399 assert!((dispatches[0].params[6] - 1.25).abs() < 1e-6);
1400 assert_eq!(dispatches[3].params[0], 1.0);
1401 assert!((dispatches[3].params[1] - 0.2).abs() < 1e-6);
1402 assert_eq!(
1403 exposure_slots(&plan[0], dispatches.len()),
1404 [Some(6), None, None, Some(0)]
1405 );
1406 assert_eq!(crate::chain::grade_data(members).len(), 6 * 16);
1407 let params = plan[0].params(Buffers::GEOMETRY);
1408 assert_eq!(params.len(), 4);
1409 for (params, dispatch) in params.iter().zip(&dispatches) {
1410 assert_eq!(params, &dispatch.params);
1411 }
1412 }
1413
1414 #[test]
1415 fn every_planned_shader_validates_without_a_gpu() {
1416 let mut chains: Vec<Chain> = finishes().into_iter().map(|(_, chain)| chain).collect();
1417 chains.extend(Style::ALL.into_iter().map(|style| style.chain()));
1418 chains.push(Chain {
1419 passes: fused_finish(0.4, 0.3),
1420 seed: 0,
1421 frame: 0,
1422 });
1423 let mut rings = Chain::new();
1424 rings.passes.push(Pass::Bloom(crate::bloom::Bloom::rings(
1425 0.5,
1426 1.0,
1427 4,
1428 crate::bloom::NEUTRAL_RADII,
1429 crate::bloom::NEUTRAL_GAINS,
1430 )));
1431 chains.push(rings);
1432 let mut seen = std::collections::BTreeSet::new();
1433 for chain in chains {
1434 for unfused in [false, true] {
1435 for pass in plan_passes(chain.passes.clone(), !unfused) {
1436 for dispatch in pass.dispatches(Buffers::GEOMETRY) {
1437 if !seen.insert(dispatch.shader.to_string()) {
1438 continue;
1439 }
1440 let module = naga::front::wgsl::parse_str(&dispatch.shader)
1441 .unwrap_or_else(|error| panic!("{}: {error}", pass.name()));
1442 naga::valid::Validator::new(
1443 naga::valid::ValidationFlags::all(),
1444 naga::valid::Capabilities::SHADER_FLOAT16_IN_FLOAT32,
1445 )
1446 .validate(&module)
1447 .unwrap_or_else(|error| panic!("{}: {error:?}", pass.name()));
1448 }
1449 }
1450 }
1451 }
1452 assert!(seen.len() > 40, "{}", seen.len());
1453 }
1454
1455 fn encode(value: f32) -> i32 {
1456 (value.clamp(0.0, 1.0) * 255.0).round() as i32
1457 }
1458
1459 fn original_point_bloom(post: &mut GpuChain, device: &wgpu::Device) {
1460 let stage = &mut post.groups[0].stages[0];
1461 *stage = make_stage(
1462 device,
1463 crate::wgsl::BLOOM_DOWN,
1464 Aux::None,
1465 stage.source,
1466 stage.target,
1467 stage.label.clone(),
1468 );
1469 for stage in &mut post.groups[0].stages[1..3] {
1470 *stage = make_stage(
1471 device,
1472 crate::wgsl::BLOOM_BLUR,
1473 Aux::None,
1474 stage.source,
1475 stage.target,
1476 stage.label.clone(),
1477 );
1478 }
1479 let stage = &mut post.groups[0].stages[3];
1480 *stage = make_stage(
1481 device,
1482 crate::wgsl::BLOOM_ADD,
1483 Aux::Bloom,
1484 stage.source,
1485 stage.target,
1486 stage.label.clone(),
1487 );
1488 }
1489
1490 #[test]
1491 #[ignore = "needs a GPU; run under pgpu"]
1492 fn styles_match_cpu_on_gpu() {
1493 let gpu = pollster::block_on(Gpu::headless()).unwrap();
1494 let source = synthetic(64, 64);
1495 let input = gpu.offscreen(64, 64, FORMAT).unwrap();
1496 let output = output_target(&gpu, 64, 64);
1497 let pixels: Vec<u16> = source
1498 .pixels
1499 .iter()
1500 .flat_map(|p| p.iter().map(|v| half(*v)))
1501 .collect();
1502 gpu.upload_rgba16(&input, &pixels).unwrap();
1503 for style in Style::ALL {
1504 let mut chain = style.chain();
1505 chain.frame = 7;
1506 chain.seed = 23;
1507 chain.passes.push(Pass::Encode);
1508 let expected = chain.apply(&source, None, None);
1509 let post = GpuChain::new(chain, &gpu.device, 64, 64);
1510 let mut encoder = gpu
1511 .device
1512 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
1513 label: Some("post test"),
1514 });
1515 post.run(&mut encoder, &input.view, None, None, &output.view, 7, 23);
1516 gpu.queue.submit(Some(encoder.finish()));
1517 let actual = gpu.readback_rgba16(&output).unwrap();
1518 for (pixel, (got, want)) in actual.chunks_exact(4).zip(&expected.pixels).enumerate() {
1519 for channel in 0..3 {
1520 let delta = (encode(float(got[channel])) - encode(want[channel])).abs();
1521 assert!(
1522 delta <= 1,
1523 "{} pixel {pixel} channel {channel}: {} vs {}",
1524 style.name(),
1525 float(got[channel]),
1526 want[channel]
1527 );
1528 }
1529 }
1530 }
1531 }
1532
1533 #[test]
1534 #[ignore = "needs a GPU; run under pgpu"]
1535 fn fused_grade_matches_unfused_at_any_warmth() {
1536 let gpu = pollster::block_on(Gpu::headless()).unwrap();
1537 let source = synthetic(64, 64);
1538 let input = gpu.offscreen(64, 64, FORMAT).unwrap();
1539 let output = output_target(&gpu, 64, 64);
1540 let pixels: Vec<u16> = source
1541 .pixels
1542 .iter()
1543 .flat_map(|p| p.iter().map(|v| half(*v)))
1544 .collect();
1545 gpu.upload_rgba16(&input, &pixels).unwrap();
1546 for warmth in [0.0, 0.5, 1.0] {
1547 let mut chain = Chain {
1548 passes: fused_finish(warmth, 0.3),
1549 seed: 23,
1550 frame: 7,
1551 };
1552 chain.passes.push(Pass::Encode);
1553 let expected = chain.apply(&source, None, None);
1554 let post = GpuChain::new(chain, &gpu.device, 64, 64);
1555 assert!(
1556 post.groups
1557 .iter()
1558 .any(|group| matches!(group.pass, Pass::Grade(_))),
1559 "warmth {warmth} did not fuse"
1560 );
1561 let mut encoder = gpu
1562 .device
1563 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
1564 label: Some("post fused test"),
1565 });
1566 post.run(&mut encoder, &input.view, None, None, &output.view, 7, 23);
1567 gpu.queue.submit(Some(encoder.finish()));
1568 let actual = gpu.readback_rgba16(&output).unwrap();
1569 for (pixel, (got, want)) in actual.chunks_exact(4).zip(&expected.pixels).enumerate() {
1570 for channel in 0..3 {
1571 let delta = (encode(float(got[channel])) - encode(want[channel])).abs();
1572 assert!(
1573 delta <= 1,
1574 "warmth {warmth} pixel {pixel} channel {channel}: {} vs {}",
1575 float(got[channel]),
1576 want[channel]
1577 );
1578 }
1579 }
1580 }
1581 }
1582
1583 #[test]
1584 #[ignore = "needs a GPU; run under pgpu"]
1585 fn staged_gpu_matches_whole_chain_without_overlay() {
1586 let gpu = pollster::block_on(Gpu::headless()).unwrap();
1587 let source = synthetic(64, 64);
1588 let input = gpu.offscreen(64, 64, FORMAT).unwrap();
1589 let middle = output_target(&gpu, 64, 64);
1590 let whole_output = output_target(&gpu, 64, 64);
1591 let staged_output = output_target(&gpu, 64, 64);
1592 let pixels: Vec<u16> = source
1593 .pixels
1594 .iter()
1595 .flat_map(|p| p.iter().map(|v| half(*v)))
1596 .collect();
1597 gpu.upload_rgba16(&input, &pixels).unwrap();
1598 let mut chain = Chain::new();
1599 chain.passes = vec![
1600 Pass::Exposure(1.0),
1601 Pass::Bloom(crate::bloom::Bloom::neutral(0.45)),
1602 Pass::Tone(crate::tone::Tone::aces()),
1603 ];
1604 let whole = GpuChain::new(chain.clone(), &gpu.device, 64, 64);
1605 let stages = GpuStages::new(&chain, &gpu.device, 64, 64);
1606 let mut encoder = gpu
1607 .device
1608 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
1609 label: Some("post staged test"),
1610 });
1611 whole.run(
1612 &mut encoder,
1613 &input.view,
1614 None,
1615 None,
1616 &whole_output.view,
1617 0,
1618 0,
1619 );
1620 stages
1621 .bloom_grade
1622 .run(&mut encoder, &input.view, None, None, &middle.view, 0, 0);
1623 stages.final_stage.run(
1624 &mut encoder,
1625 &middle.view,
1626 None,
1627 None,
1628 &staged_output.view,
1629 0,
1630 0,
1631 );
1632 gpu.queue.submit(Some(encoder.finish()));
1633 let whole_pixels = gpu.readback_rgba16(&whole_output).unwrap();
1634 let staged_pixels = gpu.readback_rgba16(&staged_output).unwrap();
1635 for (a, b) in whole_pixels.iter().zip(staged_pixels.iter()) {
1636 assert!((encode(float(*a)) - encode(float(*b))).abs() <= 1);
1637 }
1638 }
1639
1640 #[test]
1641 #[ignore = "needs a GPU; run under pgpu"]
1642 fn box_bloom_matches_the_point_gaussian() {
1643 let gpu = pollster::block_on(Gpu::headless()).unwrap();
1644 let (width, height) = (256, 128);
1645 let source = Image::from_fn(width, height, |x, y| {
1646 let a = (x as f32 - 61.0).hypot(y as f32 - 37.0);
1647 let b = (x as f32 - 194.0).hypot(y as f32 - 82.0);
1648 let light = (2.0 - a * 0.08).max(0.0) + (1.6 - b * 0.06).max(0.0);
1649 [light, light * 0.7, light * 0.4, 1.0]
1650 });
1651 let input = gpu.offscreen(width, height, FORMAT).unwrap();
1652 let original_output = output_target(&gpu, width, height);
1653 let optimized_output = output_target(&gpu, width, height);
1654 let pixels: Vec<u16> = source
1655 .pixels
1656 .iter()
1657 .flat_map(|p| p.iter().map(|v| half(*v)))
1658 .collect();
1659 gpu.upload_rgba16(&input, &pixels).unwrap();
1660 let mut chain = Chain::new();
1661 chain.passes.push(Pass::Bloom(crate::bloom::Bloom::pyramid(
1662 0.45, 1.0, 0.0, 0.0, 0.0, 4, 16.0, 2,
1663 )));
1664 let expected = chain.apply(&source, None, None);
1665 let mut original = GpuChain::new(chain.clone(), &gpu.device, width, height);
1666 original_point_bloom(&mut original, &gpu.device);
1667 let optimized = GpuChain::new(chain, &gpu.device, width, height);
1668 let mut encoder = gpu
1669 .device
1670 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
1671 label: Some("bloom comparison"),
1672 });
1673 original.run(
1674 &mut encoder,
1675 &input.view,
1676 None,
1677 None,
1678 &original_output.view,
1679 0,
1680 0,
1681 );
1682 optimized.run(
1683 &mut encoder,
1684 &input.view,
1685 None,
1686 None,
1687 &optimized_output.view,
1688 0,
1689 0,
1690 );
1691 gpu.queue.submit(Some(encoder.finish()));
1692 let old = gpu.readback_rgba16(&original_output).unwrap();
1693 let new = gpu.readback_rgba16(&optimized_output).unwrap();
1694 for (actual, want) in new.chunks_exact(4).zip(&expected.pixels) {
1695 for channel in 0..3 {
1696 assert!((encode(float(actual[channel])) - encode(want[channel])).abs() <= 2);
1697 }
1698 }
1699 let mut differences: Vec<i32> = old
1700 .chunks_exact(4)
1701 .zip(new.chunks_exact(4))
1702 .flat_map(|(a, b)| {
1703 (0..3).map(move |channel| {
1704 (encode(float(a[channel])) - encode(float(b[channel]))).abs()
1705 })
1706 })
1707 .collect();
1708 differences.sort_unstable();
1709 let median = differences[differences.len() / 2];
1710 let maximum = differences.last().copied().unwrap();
1711 assert_eq!(median, 0);
1712 assert!(maximum <= 2);
1713 let mut either = 0;
1714 let mut both = 0;
1715 for ((base, a), b) in pixels
1716 .chunks_exact(4)
1717 .zip(old.chunks_exact(4))
1718 .zip(new.chunks_exact(4))
1719 {
1720 let old_bright = float(a[0]) - float(base[0]) > 1.0 / 255.0;
1721 let new_bright = float(b[0]) - float(base[0]) > 1.0 / 255.0;
1722 if old_bright || new_bright {
1723 either += 1;
1724 if old_bright && new_bright {
1725 both += 1;
1726 }
1727 }
1728 }
1729 assert!(either > 0);
1730 let overlap = both as f32 / either as f32;
1731 println!("box bloom median {median}/255 max {maximum}/255 footprint {overlap:.3}");
1732 assert!(overlap >= 0.985);
1733 }
1734
1735 #[test]
1736 #[ignore = "needs a GPU; run under pgpu"]
1737 fn geometry_views_match_cpu() {
1738 let gpu = pollster::block_on(Gpu::headless()).unwrap();
1739 let source = synthetic(64, 64);
1740 let input = gpu.offscreen(64, 64, FORMAT).unwrap();
1741 let output = output_target(&gpu, 64, 64);
1742 let pixels: Vec<u16> = source
1743 .pixels
1744 .iter()
1745 .flat_map(|p| p.iter().map(|v| half(*v)))
1746 .collect();
1747 gpu.upload_rgba16(&input, &pixels).unwrap();
1748 let depth_values = vec![0.5; 4096];
1749 let normals: Vec<[f32; 3]> = (0..4096)
1750 .map(|i| {
1751 if i % 64 < 32 {
1752 [0.0, 0.0, 1.0]
1753 } else {
1754 [0.0, 0.5, 0.5]
1755 }
1756 })
1757 .collect();
1758 let depth_texture = gpu.device.create_texture(&wgpu::TextureDescriptor {
1759 label: Some("post test depth"),
1760 size: wgpu::Extent3d {
1761 width: 64,
1762 height: 64,
1763 depth_or_array_layers: 1,
1764 },
1765 mip_level_count: 1,
1766 sample_count: 1,
1767 dimension: wgpu::TextureDimension::D2,
1768 format: wgpu::TextureFormat::Depth32Float,
1769 usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::RENDER_ATTACHMENT,
1770 view_formats: &[],
1771 });
1772 let depth_view = depth_texture.create_view(&wgpu::TextureViewDescriptor::default());
1773 let mut depth_encoder =
1774 gpu.device
1775 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
1776 label: Some("post depth clear"),
1777 });
1778 {
1779 let _pass = depth_encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
1780 multiview_mask: None,
1781 label: Some("post depth clear"),
1782 color_attachments: &[],
1783 depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
1784 view: &depth_view,
1785 depth_ops: Some(wgpu::Operations {
1786 load: wgpu::LoadOp::Clear(0.5),
1787 store: wgpu::StoreOp::Store,
1788 }),
1789 stencil_ops: None,
1790 }),
1791 timestamp_writes: None,
1792 occlusion_query_set: None,
1793 });
1794 }
1795 gpu.queue.submit(Some(depth_encoder.finish()));
1796 let normal_target = gpu.offscreen(64, 64, FORMAT).unwrap();
1797 let normal_pixels: Vec<u16> = normals
1798 .iter()
1799 .flat_map(|n| [half(n[0]), half(n[1]), half(n[2]), half(1.0)])
1800 .collect();
1801 gpu.upload_rgba16(&normal_target, &normal_pixels).unwrap();
1802 for style in [
1803 Style::Cel,
1804 Style::RubberHose,
1805 Style::Comic,
1806 Style::ModernComic,
1807 Style::TiltShift,
1808 ] {
1809 let mut chain = style.chain();
1810 chain.passes.push(Pass::Encode);
1811 let expected = chain.apply(&source, Some(&depth_values), Some(&normals));
1812 let post = GpuChain::new(chain, &gpu.device, 64, 64);
1813 let mut encoder = gpu
1814 .device
1815 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
1816 label: Some("post geometry test"),
1817 });
1818 post.run(
1819 &mut encoder,
1820 &input.view,
1821 Some(&depth_view),
1822 Some(&normal_target.view),
1823 &output.view,
1824 0,
1825 0,
1826 );
1827 gpu.queue.submit(Some(encoder.finish()));
1828 let actual = gpu.readback_rgba16(&output).unwrap();
1829 for (pixel, (got, want)) in actual.chunks_exact(4).zip(&expected.pixels).enumerate() {
1830 for channel in 0..3 {
1831 let delta = (encode(float(got[channel])) - encode(want[channel])).abs();
1832 assert!(
1833 delta <= 1,
1834 "{} pixel {pixel} channel {channel}: {} vs {}",
1835 style.name(),
1836 float(got[channel]),
1837 want[channel]
1838 );
1839 }
1840 }
1841 }
1842 }
1843
1844 #[test]
1845 #[ignore = "needs a GPU; run under pgpu"]
1846 fn pipeline_creation_and_resize() {
1847 let gpu = pollster::block_on(Gpu::headless()).unwrap();
1848 for style in Style::ALL {
1849 let mut post = GpuChain::new(style.chain(), &gpu.device, 64, 64);
1850 assert_eq!(post.size(), (64, 64));
1851 post.resize(113, 57);
1852 assert_eq!(post.size(), (113, 57));
1853 for group in &post.groups {
1854 if let Some(low) = &group.low {
1855 assert_eq!(
1856 (low.a._texture.width(), low.a._texture.height()),
1857 ((113 / group.factor).max(1), (57 / group.factor).max(1))
1858 );
1859 assert_eq!(low.b._texture.size(), low.a._texture.size());
1860 }
1861 }
1862 }
1863 }
1864
1865 #[test]
1866 #[ignore = "needs a GPU; run under pgpu"]
1867 fn benchmark_4k_styles() {
1868 let gpu = pollster::block_on(Gpu::headless()).unwrap();
1869 let input = gpu.offscreen(3840, 2160, FORMAT).unwrap();
1870 let output = output_target(&gpu, 3840, 2160);
1871 let normal = gpu.offscreen(3840, 2160, FORMAT).unwrap();
1872 let depth = gpu.device.create_texture(&wgpu::TextureDescriptor {
1873 label: Some("post benchmark depth"),
1874 size: wgpu::Extent3d {
1875 width: 3840,
1876 height: 2160,
1877 depth_or_array_layers: 1,
1878 },
1879 mip_level_count: 1,
1880 sample_count: 1,
1881 dimension: wgpu::TextureDimension::D2,
1882 format: wgpu::TextureFormat::Depth32Float,
1883 usage: wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::TEXTURE_BINDING,
1884 view_formats: &[],
1885 });
1886 let depth_view = depth.create_view(&Default::default());
1887 let mut clear = gpu
1888 .device
1889 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
1890 label: Some("post benchmark geometry clear"),
1891 });
1892 {
1893 let _pass = clear.begin_render_pass(&wgpu::RenderPassDescriptor {
1894 multiview_mask: None,
1895 label: Some("post benchmark geometry"),
1896 color_attachments: &[Some(wgpu::RenderPassColorAttachment {
1897 view: &normal.view,
1898 resolve_target: None,
1899 depth_slice: None,
1900 ops: wgpu::Operations {
1901 load: wgpu::LoadOp::Clear(wgpu::Color {
1902 r: 0.0,
1903 g: 0.0,
1904 b: 1.0,
1905 a: 1.0,
1906 }),
1907 store: wgpu::StoreOp::Store,
1908 },
1909 })],
1910 depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
1911 view: &depth_view,
1912 depth_ops: Some(wgpu::Operations {
1913 load: wgpu::LoadOp::Clear(0.5),
1914 store: wgpu::StoreOp::Store,
1915 }),
1916 stencil_ops: None,
1917 }),
1918 timestamp_writes: None,
1919 occlusion_query_set: None,
1920 });
1921 }
1922 gpu.queue.submit(Some(clear.finish()));
1923 let mut timer = pfx_gpu::GpuProfiler::new(&gpu.device, &gpu.queue);
1924 let mut over_budget = Vec::new();
1925 let mut turns = pfx_gpu::pace::Turns::default();
1926 for style in Style::ALL {
1927 let post = turns.cpu(|| GpuChain::new(style.chain(), &gpu.device, 3840, 2160));
1928 let mut samples = Vec::new();
1929 for attempt in 0..12 {
1930 let mut encoder =
1931 gpu.device
1932 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
1933 label: Some("post benchmark"),
1934 });
1935 let geometry = matches!(
1936 style,
1937 Style::Cel | Style::RubberHose | Style::Comic | Style::ModernComic
1938 );
1939 post.run_timed(
1940 &mut encoder,
1941 &input.view,
1942 geometry.then_some(&depth_view),
1943 geometry.then_some(&normal.view),
1944 &output.view,
1945 7,
1946 23,
1947 Some(&mut timer),
1948 );
1949 let slot = timer.finish(&mut encoder);
1950 gpu.queue.submit(Some(encoder.finish()));
1951 if let Some(slot) = slot {
1952 timer.submitted(slot);
1953 }
1954 gpu.device
1955 .poll(wgpu::PollType::wait_indefinitely())
1956 .unwrap();
1957 let timings = timer.collect(&gpu.device);
1958 if attempt == 11
1959 && matches!(
1960 style,
1961 Style::Cel | Style::RubberHose | Style::Comic | Style::ModernComic
1962 )
1963 {
1964 for pass in timings.iter().flatten() {
1965 println!(
1966 "{} {} {:.3} ms",
1967 style.name(),
1968 pass.label,
1969 pass.milliseconds
1970 );
1971 }
1972 }
1973 let total: f64 = timings.iter().flatten().map(|pass| pass.milliseconds).sum();
1974 turns.add(total);
1975 if attempt > 0 {
1976 samples.push(total);
1977 }
1978 }
1979 samples.sort_by(f64::total_cmp);
1980 println!("{} {:.3} ms", style.name(), samples[5]);
1981 if matches!(
1982 style,
1983 Style::Cel | Style::RubberHose | Style::Comic | Style::ModernComic
1984 ) && samples[5] >= 2.0
1985 {
1986 over_budget.push(format!("{} {:.3} ms", style.name(), samples[5]));
1987 }
1988 }
1989 assert!(over_budget.is_empty(), "{}", over_budget.join(", "));
1990 }
1991
1992 #[test]
1993 #[ignore = "needs a GPU; run under pgpu"]
1994 fn normals_reach_rim_chain() {
1995 let gpu = pollster::block_on(Gpu::headless()).unwrap();
1996 let input = gpu.offscreen(64, 64, FORMAT).unwrap();
1997 let normal = gpu.offscreen(64, 64, FORMAT).unwrap();
1998 let output = output_target(&gpu, 64, 64);
1999 gpu.upload_rgba16(&input, &vec![half(0.2); 64 * 64 * 4])
2000 .unwrap();
2001 let normals: Vec<u16> = (0..64 * 64)
2002 .flat_map(|i| {
2003 let angled = i % 64 >= 16 && i % 64 < 48;
2004 [
2005 half(if angled { 0.98 } else { 0.0 }),
2006 half(0.0),
2007 half(if angled { 0.2 } else { 1.0 }),
2008 half(1.0),
2009 ]
2010 })
2011 .collect();
2012 gpu.upload_rgba16(&normal, &normals).unwrap();
2013 let chain = GpuChain::new(
2014 Chain {
2015 passes: vec![Pass::Rim(crate::fx::Rim {
2016 strength: 0.8,
2017 width: 0.3,
2018 color: [0.3, 0.6, 1.0],
2019 })],
2020 frame: 0,
2021 seed: 0,
2022 },
2023 &gpu.device,
2024 64,
2025 64,
2026 );
2027 let count = |normal_view: Option<&wgpu::TextureView>| {
2028 let mut encoder = gpu
2029 .device
2030 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
2031 label: Some("rim normal test"),
2032 });
2033 chain.run(
2034 &mut encoder,
2035 &input.view,
2036 None,
2037 normal_view,
2038 &output.view,
2039 0,
2040 0,
2041 );
2042 gpu.queue.submit(Some(encoder.finish()));
2043 gpu.readback_rgba16(&output)
2044 .unwrap()
2045 .chunks_exact(4)
2046 .filter(|pixel| float(pixel[2]) > 0.25)
2047 .count()
2048 };
2049 let without = count(None);
2050 let with = count(Some(&normal.view));
2051 assert!(with > without + 1000, "{without} vs {with}");
2052 }
2053
2054 #[test]
2055 #[ignore = "needs a GPU; run under pgpu"]
2056 fn benchmark_bloom_sizes() {
2057 let gpu = pollster::block_on(Gpu::headless()).unwrap();
2058 let mut timer = pfx_gpu::GpuProfiler::new(&gpu.device, &gpu.queue);
2059 let mut turns = pfx_gpu::pace::Turns::default();
2060 for (width, height) in [(1586, 992), (3840, 2160)] {
2061 let input = gpu.offscreen(width, height, FORMAT).unwrap();
2062 let output = output_target(&gpu, width, height);
2063 for (name, bloom) in [
2064 ("box", crate::bloom::Bloom::neutral(0.45)),
2065 ("tent", crate::bloom::Bloom::tent(0.2, 1.0, 4, 16.0, 2)),
2066 (
2067 "rings",
2068 crate::bloom::Bloom::rings(
2069 0.5,
2070 1.0,
2071 4,
2072 crate::bloom::NEUTRAL_RADII,
2073 crate::bloom::NEUTRAL_GAINS,
2074 ),
2075 ),
2076 ] {
2077 let mut chain = Chain::new();
2078 chain.passes.push(Pass::Bloom(bloom));
2079 for version in ["before", "after"] {
2080 let mut post =
2081 turns.cpu(|| GpuChain::new(chain.clone(), &gpu.device, width, height));
2082 if version == "before" && name == "box" {
2083 original_point_bloom(&mut post, &gpu.device);
2084 }
2085 let mut samples = Vec::new();
2086 for attempt in 0..15 {
2087 let mut encoder =
2088 gpu.device
2089 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
2090 label: Some("bloom benchmark"),
2091 });
2092 post.run_timed(
2093 &mut encoder,
2094 &input.view,
2095 None,
2096 None,
2097 &output.view,
2098 0,
2099 0,
2100 Some(&mut timer),
2101 );
2102 let slot = timer.finish(&mut encoder);
2103 gpu.queue.submit(Some(encoder.finish()));
2104 if let Some(slot) = slot {
2105 timer.submitted(slot);
2106 }
2107 gpu.device
2108 .poll(wgpu::PollType::wait_indefinitely())
2109 .unwrap();
2110 let timings = timer.collect(&gpu.device);
2111 let total: f64 =
2112 timings.iter().flatten().map(|pass| pass.milliseconds).sum();
2113 turns.add(total);
2114 if attempt > 0 {
2115 samples.push(total);
2116 }
2117 }
2118 samples.sort_by(f64::total_cmp);
2119 println!("{name} {width}x{height} {version} {:.3} ms", samples[7]);
2120 }
2121 }
2122 }
2123 }
2124
2125 fn frame_finish() -> Chain {
2126 Chain {
2127 passes: vec![
2128 Pass::Exposure(1.25),
2129 Pass::Bloom(crate::bloom::Bloom::neutral(0.2)),
2130 Pass::Vignette(crate::grade::Vignette::power(0.4, 2.0, 1.0, 1.0)),
2131 Pass::Saturation(crate::grade::Saturation::rec709(1.1)),
2132 Pass::Warmth(crate::grade::Warmth::linear(0.0)),
2133 Pass::Tone(crate::tone::Tone::aces()),
2134 ],
2135 seed: 0,
2136 frame: 0,
2137 }
2138 }
2139
2140 fn as_live(mut chain: Chain) -> Chain {
2141 chain.passes.insert(0, Pass::Exposure(1.0));
2142 chain.passes.push(Pass::Encode);
2143 chain
2144 }
2145
2146 fn finishes() -> Vec<(&'static str, Chain)> {
2147 let mut out = vec![("frame finish", as_live(frame_finish()))];
2148 out.extend(
2149 Style::ALL
2150 .into_iter()
2151 .map(|style| (style.name(), as_live(style.chain()))),
2152 );
2153 out
2154 }
2155
2156 fn uses_geometry(chain: &Chain) -> bool {
2157 chain.passes.iter().any(|pass| {
2158 matches!(
2159 pass,
2160 Pass::Cel(_)
2161 | Pass::Outline(_)
2162 | Pass::Cavity(_)
2163 | Pass::Rim(_)
2164 | Pass::Comic(_)
2165 | Pass::RubberHose
2166 | Pass::ModernComic(_)
2167 | Pass::TiltShift(_)
2168 )
2169 })
2170 }
2171
2172 fn captured_frame(width: u32, height: u32) -> Option<Image> {
2173 let path = std::env::var_os("POST_FRAME_PNG")?;
2174 let bytes = std::fs::read(path).ok()?;
2175 let mut decoder = png::Decoder::new(std::io::Cursor::new(bytes));
2176 decoder.set_transformations(png::Transformations::EXPAND);
2177 let mut reader = decoder.read_info().ok()?;
2178 let mut buf = vec![0u8; reader.output_buffer_size()?];
2179 let info = reader.next_frame(&mut buf).ok()?;
2180 if (info.width, info.height) != (width, height) || info.color_type != png::ColorType::Rgba {
2181 return None;
2182 }
2183 Some(Image::from_fn(width, height, |x, y| {
2184 let at = ((y * width + x) * 4) as usize;
2185 let channel = |c: usize| pfx_materials::linear_channel(f32::from(buf[at + c]) / 255.0);
2186 [channel(0), channel(1), channel(2), 1.0]
2187 }))
2188 }
2189
2190 fn lit_scene(width: u32, height: u32) -> Image {
2191 let lights = [(0.18, 0.22, 6.0), (0.71, 0.35, 3.5), (0.44, 0.8, 9.0)];
2192 Image::from_fn(width, height, |x, y| {
2193 let u = x as f32 / width as f32;
2194 let v = y as f32 / height as f32;
2195 let mut c = [0.05 + 0.4 * u, 0.08 + 0.3 * v, 0.12 + 0.2 * (u * v), 1.0];
2196 for (lx, ly, power) in lights {
2197 let d = ((u - lx) * 16.0).hypot((v - ly) * 9.0);
2198 let glow = power / (1.0 + d * d * 40.0);
2199 c[0] += glow;
2200 c[1] += glow * 0.8;
2201 c[2] += glow * 0.55;
2202 }
2203 let stripe = ((x / 37 + y / 23) % 5) as f32 * 0.04;
2204 [c[0] + stripe, c[1] + stripe, c[2], 1.0]
2205 })
2206 }
2207
2208 fn source_4k(gpu: &Gpu) -> (OffscreenTarget, &'static str) {
2209 let (width, height) = (3840, 2160);
2210 let (image, name) = match captured_frame(width, height) {
2211 Some(image) => (image, "a captured frame"),
2212 None => (lit_scene(width, height), "a lit synthetic scene"),
2213 };
2214 let input = gpu.offscreen(width, height, FORMAT).unwrap();
2215 let pixels: Vec<u16> = image
2216 .pixels
2217 .iter()
2218 .flat_map(|p| p.iter().map(|v| half(*v)))
2219 .collect();
2220 gpu.upload_rgba16(&input, &pixels).unwrap();
2221 (input, name)
2222 }
2223
2224 struct Geometry {
2225 _depth: wgpu::Texture,
2226 depth: wgpu::TextureView,
2227 normal: OffscreenTarget,
2228 }
2229
2230 fn geometry_4k(gpu: &Gpu) -> Geometry {
2231 let normal = gpu.offscreen(3840, 2160, FORMAT).unwrap();
2232 let depth = gpu.device.create_texture(&wgpu::TextureDescriptor {
2233 label: Some("post finish depth"),
2234 size: wgpu::Extent3d {
2235 width: 3840,
2236 height: 2160,
2237 depth_or_array_layers: 1,
2238 },
2239 mip_level_count: 1,
2240 sample_count: 1,
2241 dimension: wgpu::TextureDimension::D2,
2242 format: wgpu::TextureFormat::Depth32Float,
2243 usage: wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::TEXTURE_BINDING,
2244 view_formats: &[],
2245 });
2246 let view = depth.create_view(&Default::default());
2247 let mut clear = gpu
2248 .device
2249 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
2250 label: Some("post finish geometry clear"),
2251 });
2252 {
2253 let _pass = clear.begin_render_pass(&wgpu::RenderPassDescriptor {
2254 multiview_mask: None,
2255 label: Some("post finish geometry"),
2256 color_attachments: &[Some(wgpu::RenderPassColorAttachment {
2257 view: &normal.view,
2258 resolve_target: None,
2259 depth_slice: None,
2260 ops: wgpu::Operations {
2261 load: wgpu::LoadOp::Clear(wgpu::Color {
2262 r: 0.0,
2263 g: 0.0,
2264 b: 1.0,
2265 a: 1.0,
2266 }),
2267 store: wgpu::StoreOp::Store,
2268 },
2269 })],
2270 depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
2271 view: &view,
2272 depth_ops: Some(wgpu::Operations {
2273 load: wgpu::LoadOp::Clear(0.5),
2274 store: wgpu::StoreOp::Store,
2275 }),
2276 stencil_ops: None,
2277 }),
2278 timestamp_writes: None,
2279 occlusion_query_set: None,
2280 });
2281 }
2282 gpu.queue.submit(Some(clear.finish()));
2283 Geometry {
2284 _depth: depth,
2285 depth: view,
2286 normal,
2287 }
2288 }
2289
2290 fn median(values: &mut [f64]) -> f64 {
2291 values.sort_by(f64::total_cmp);
2292 values[values.len() / 2]
2293 }
2294
2295 #[test]
2296 #[ignore = "needs a GPU; run under pgpu"]
2297 fn fused_matches_unfused_at_4k() {
2298 let gpu = pollster::block_on(Gpu::headless()).unwrap();
2299 let mut turns = pfx_gpu::pace::Turns::default();
2300 let (input, source) = turns.cpu(|| source_4k(&gpu));
2301 let geometry = turns.cpu(|| geometry_4k(&gpu));
2302 let fused_output = output_target(&gpu, 3840, 2160);
2303 let reference_output = output_target(&gpu, 3840, 2160);
2304 let levels: Vec<i16> = turns.cpu(|| {
2305 (0..=u16::MAX)
2306 .map(|bits| encode(float(bits)) as i16)
2307 .collect()
2308 });
2309 let mut failures = Vec::new();
2310 println!("3840x2160 from {source}");
2311 for (name, chain) in finishes() {
2312 let geometric = uses_geometry(&chain);
2313 let fused = turns.cpu(|| GpuChain::new(chain.clone(), &gpu.device, 3840, 2160));
2314 let reference = turns.cpu(|| GpuChain::unfused(chain, &gpu.device, 3840, 2160));
2315 let mut worst = 0;
2316 let mut differing = 0usize;
2317 let exposures: &[f32] = if name == "frame finish" {
2318 &[1.0, 0.8]
2319 } else {
2320 &[0.8]
2321 };
2322 for &exposure in exposures {
2323 for (post, output) in [(&fused, &fused_output), (&reference, &reference_output)] {
2324 let mut encoder =
2325 gpu.device
2326 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
2327 label: Some("post fused comparison"),
2328 });
2329 post.run_timed_exposed(
2330 &mut encoder,
2331 &input.view,
2332 geometric.then_some(&geometry.depth),
2333 geometric.then_some(&geometry.normal.view),
2334 &output.view,
2335 7,
2336 23,
2337 exposure,
2338 None,
2339 );
2340 turns.time(|| {
2341 gpu.queue.submit(Some(encoder.finish()));
2342 gpu.device
2343 .poll(wgpu::PollType::wait_indefinitely())
2344 .unwrap();
2345 });
2346 }
2347 let got = turns.cpu(|| gpu.readback_rgba16(&fused_output).unwrap());
2348 let want = turns.cpu(|| gpu.readback_rgba16(&reference_output).unwrap());
2349 for (pixel, (a, b)) in got.chunks_exact(4).zip(want.chunks_exact(4)).enumerate() {
2350 if pixel % 262_144 == 0 {
2351 turns.poll();
2352 }
2353 if a == b {
2354 continue;
2355 }
2356 for channel in 0..3 {
2357 let delta = (levels[usize::from(a[channel])]
2358 - levels[usize::from(b[channel])])
2359 .abs();
2360 if delta > 0 {
2361 differing += 1;
2362 worst = worst.max(delta);
2363 }
2364 }
2365 }
2366 }
2367 println!(
2368 "{name}: {} passes fused into {}, worst {worst}/255, {differing} channels differ",
2369 reference
2370 .groups
2371 .iter()
2372 .map(|group| group.stages.len())
2373 .sum::<usize>(),
2374 fused
2375 .groups
2376 .iter()
2377 .map(|group| group.stages.len())
2378 .sum::<usize>(),
2379 );
2380 if worst > 1 {
2381 failures.push(format!("{name} {worst}/255"));
2382 }
2383 }
2384 assert!(failures.is_empty(), "{}", failures.join(", "));
2385 }
2386
2387 #[test]
2388 #[ignore = "needs a GPU; run under pgpu"]
2389 fn benchmark_4k_finish_breakdown() {
2390 let gpu = pollster::block_on(Gpu::headless()).unwrap();
2391 let mut turns = pfx_gpu::pace::Turns::default();
2392 let (input, source) = turns.cpu(|| source_4k(&gpu));
2393 let output = output_target(&gpu, 3840, 2160);
2394 let geometry = turns.cpu(|| geometry_4k(&gpu));
2395 let mut timer = pfx_gpu::GpuProfiler::new(&gpu.device, &gpu.queue);
2396 let frames = std::env::var("POST_BENCH_FRAMES")
2397 .ok()
2398 .and_then(|value| value.parse::<usize>().ok())
2399 .unwrap_or(300);
2400 println!("3840x2160 from {source}, medians over {frames} frames");
2401 let mut frame = f64::INFINITY;
2402 for (name, chain) in finishes() {
2403 let geometric = uses_geometry(&chain);
2404 let post = turns.cpu(|| GpuChain::new(chain, &gpu.device, 3840, 2160));
2405 let mut totals = Vec::with_capacity(frames);
2406 let mut passes: Vec<(String, Vec<f64>)> = Vec::new();
2407 for attempt in 0..frames + 5 {
2408 let mut encoder =
2409 gpu.device
2410 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
2411 label: Some("post finish benchmark"),
2412 });
2413 post.run_timed_exposed(
2414 &mut encoder,
2415 &input.view,
2416 geometric.then_some(&geometry.depth),
2417 geometric.then_some(&geometry.normal.view),
2418 &output.view,
2419 attempt as u32,
2420 23,
2421 1.0,
2422 Some(&mut timer),
2423 );
2424 let slot = timer.finish(&mut encoder);
2425 gpu.queue.submit(Some(encoder.finish()));
2426 if let Some(slot) = slot {
2427 timer.submitted(slot);
2428 }
2429 gpu.device
2430 .poll(wgpu::PollType::wait_indefinitely())
2431 .unwrap();
2432 let timings = timer.collect(&gpu.device);
2433 let total: f64 = timings.iter().flatten().map(|pass| pass.milliseconds).sum();
2434 turns.add(total);
2435 if attempt < 5 {
2436 continue;
2437 }
2438 totals.push(total);
2439 for pass in timings.iter().flatten() {
2440 let label = pass
2441 .label
2442 .strip_prefix("post ")
2443 .unwrap_or(&pass.label)
2444 .to_owned();
2445 match passes.iter_mut().find(|(known, _)| *known == label) {
2446 Some((_, values)) => values.push(pass.milliseconds),
2447 None => passes.push((label, vec![pass.milliseconds])),
2448 }
2449 }
2450 }
2451 let total = median(&mut totals);
2452 if name == "frame finish" {
2453 frame = total;
2454 }
2455 let breakdown: Vec<String> = passes
2456 .iter_mut()
2457 .map(|(label, values)| format!("{label} {:.3}", median(values)))
2458 .collect();
2459 println!("{name} {total:.3} ms: {}", breakdown.join(", "));
2460 }
2461 println!("turns taken {}", turns.taken());
2462 assert!(frame.is_finite());
2463 }
2464
2465 fn region_chains() -> Vec<(&'static str, Chain)> {
2466 let finish = |extra: Vec<Pass>| {
2467 let mut passes = vec![Pass::Exposure(1.1)];
2468 passes.extend(extra);
2469 passes.push(Pass::Vignette(crate::grade::Vignette::power(
2470 0.4, 2.0, 1.0, 1.0,
2471 )));
2472 passes.push(Pass::Tone(crate::tone::Tone::aces()));
2473 passes.push(Pass::Grain(crate::grade::Grain::clamped(0.03, 0.5)));
2474 passes.push(Pass::Encode);
2475 Chain {
2476 passes,
2477 seed: 31,
2478 frame: 5,
2479 }
2480 };
2481 vec![
2482 ("pointwise finish", finish(vec![])),
2483 ("kuwahara", finish(vec![Pass::Kuwahara(3)])),
2484 (
2485 "watercolor",
2486 finish(vec![Pass::Watercolor(crate::fx::Watercolor {
2487 darkening: 0.5,
2488 grain: 0.2,
2489 scale: 6.0,
2490 })]),
2491 ),
2492 (
2493 "weave and aberration",
2494 finish(vec![Pass::Weave(3.0), Pass::Aberration(2.5)]),
2495 ),
2496 (
2497 "two kernels",
2498 finish(vec![Pass::Kuwahara(2), Pass::Sepia(0.5), Pass::Kuwahara(4)]),
2499 ),
2500 ]
2501 }
2502
2503 fn upload_source(gpu: &Gpu, source: &Image) -> OffscreenTarget {
2504 let input = gpu.offscreen(source.width, source.height, FORMAT).unwrap();
2505 let pixels: Vec<u16> = source
2506 .pixels
2507 .iter()
2508 .flat_map(|p| p.iter().map(|v| half(*v)))
2509 .collect();
2510 gpu.upload_rgba16(&input, &pixels).unwrap();
2511 input
2512 }
2513
2514 fn post_into(
2515 gpu: &Gpu,
2516 post: &GpuChain,
2517 input: &OffscreenTarget,
2518 output: &OffscreenTarget,
2519 chain: &Chain,
2520 rects: Option<&[[u32; 4]]>,
2521 ) -> Option<bool> {
2522 let mut encoder = gpu
2523 .device
2524 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
2525 label: Some("post region test"),
2526 });
2527 let scissored = match rects {
2528 Some(rects) => Some(post.run_region(
2529 &mut encoder,
2530 &input.view,
2531 None,
2532 None,
2533 &output.view,
2534 chain.frame,
2535 chain.seed,
2536 1.0,
2537 rects,
2538 None,
2539 )),
2540 None => {
2541 post.run(
2542 &mut encoder,
2543 &input.view,
2544 None,
2545 None,
2546 &output.view,
2547 chain.frame,
2548 chain.seed,
2549 );
2550 None
2551 }
2552 };
2553 gpu.queue.submit(Some(encoder.finish()));
2554 scissored
2555 }
2556
2557 #[test]
2558 #[ignore = "needs a GPU; run under pgpu"]
2559 fn region_post_equals_the_full_post_inside_its_rects() {
2560 let gpu = pollster::block_on(Gpu::headless()).unwrap();
2561 let (width, height) = (203u32, 157u32);
2562 let source = synthetic(width, height);
2563 let input = upload_source(&gpu, &source);
2564 let sets: [&[[u32; 4]]; 4] = [
2565 &[[13, 7, 61, 41]],
2566 &[[0, 0, 9, 9], [100, 60, 190, 140], [150, 3, 203, 20]],
2567 &[[190, 140, 400, 400]],
2568 &[[20, 20, 140, 120], [60, 40, 180, 150]],
2569 ];
2570 for (name, chain) in region_chains() {
2571 let post = GpuChain::new(chain.clone(), &gpu.device, width, height);
2572 assert!(post.region_reach().is_some(), "{name}");
2573 let full = output_target(&gpu, width, height);
2574 post_into(&gpu, &post, &input, &full, &chain, None);
2575 let want = gpu.readback_rgba16(&full).unwrap();
2576 for rects in sets {
2577 let region = output_target(&gpu, width, height);
2578 let scissored = post_into(&gpu, &post, &input, ®ion, &chain, Some(rects));
2579 assert_eq!(scissored, Some(true), "{name}");
2580 let got = gpu.readback_rgba16(®ion).unwrap();
2581 let mut inside_pixels = 0;
2582 for y in 0..height {
2583 for x in 0..width {
2584 let inside = rects
2585 .iter()
2586 .any(|r| x >= r[0] && x < r[2] && y >= r[1] && y < r[3]);
2587 let at = ((y * width + x) * 4) as usize;
2588 if inside {
2589 inside_pixels += 1;
2590 for channel in 0..4 {
2591 let delta =
2592 (float(got[at + channel]) - float(want[at + channel])).abs();
2593 assert!(
2594 delta <= 1.0 / 255.0,
2595 "{name} {rects:?} ({x},{y}) channel {channel}: {} vs {}",
2596 float(got[at + channel]),
2597 float(want[at + channel])
2598 );
2599 }
2600 } else {
2601 assert!(
2602 got[at..at + 4].iter().all(|v| *v == 0),
2603 "{name} {rects:?} ({x},{y}) outside the rects was written"
2604 );
2605 }
2606 }
2607 }
2608 assert!(inside_pixels > 0);
2609 }
2610 }
2611 }
2612
2613 #[test]
2614 #[ignore = "needs a GPU; run under pgpu"]
2615 fn region_post_falls_back_to_the_full_post_for_whole_frame_kernels() {
2616 let gpu = pollster::block_on(Gpu::headless()).unwrap();
2617 let (width, height) = (96u32, 64u32);
2618 let source = synthetic(width, height);
2619 let input = upload_source(&gpu, &source);
2620 let chains = [
2621 as_live(frame_finish()),
2622 Chain {
2623 passes: vec![Pass::Exposure(1.0), Pass::Distortion(0.05), Pass::Encode],
2624 seed: 3,
2625 frame: 2,
2626 },
2627 Chain {
2628 passes: vec![
2629 Pass::Exposure(1.0),
2630 Pass::Pixel(crate::fx::Pixelate {
2631 size: 5,
2632 levels: 0,
2633 palette: false,
2634 }),
2635 Pass::Encode,
2636 ],
2637 seed: 3,
2638 frame: 2,
2639 },
2640 ];
2641 for chain in chains {
2642 let post = GpuChain::new(chain.clone(), &gpu.device, width, height);
2643 assert_eq!(post.region_reach(), None);
2644 let full = output_target(&gpu, width, height);
2645 post_into(&gpu, &post, &input, &full, &chain, None);
2646 let region = output_target(&gpu, width, height);
2647 let scissored = post_into(
2648 &gpu,
2649 &post,
2650 &input,
2651 ®ion,
2652 &chain,
2653 Some(&[[10, 10, 30, 30]]),
2654 );
2655 assert_eq!(scissored, Some(false));
2656 assert_eq!(
2657 gpu.readback_rgba16(&full).unwrap(),
2658 gpu.readback_rgba16(®ion).unwrap()
2659 );
2660 }
2661 }
2662
2663 #[test]
2664 #[ignore = "needs a GPU; run under pgpu"]
2665 fn benchmark_region_post_4k() {
2666 let gpu = pollster::block_on(Gpu::headless()).unwrap();
2667 let mut timer = pfx_gpu::GpuProfiler::new(&gpu.device, &gpu.queue);
2668 let mut turns = pfx_gpu::pace::Turns::default();
2669 let (width, height) = (3840u32, 2160u32);
2670 let input = gpu.offscreen(width, height, FORMAT).unwrap();
2671 let output = output_target(&gpu, width, height);
2672 let mut chains = vec![(
2673 "standard",
2674 Chain {
2675 passes: vec![
2676 Pass::Exposure(1.0),
2677 Pass::Tone(crate::tone::Tone::aces()),
2678 Pass::Encode,
2679 ],
2680 seed: 0,
2681 frame: 0,
2682 },
2683 )];
2684 chains.extend(
2685 region_chains().into_iter().filter(|(name, _)| {
2686 matches!(*name, "pointwise finish" | "kuwahara" | "two kernels")
2687 }),
2688 );
2689 let rects: [(&str, [u32; 4]); 2] = [
2690 ("64x64", [1200, 800, 1264, 864]),
2691 ("512x256", [1200, 800, 1712, 1056]),
2692 ];
2693 for (name, chain) in chains {
2694 let post = turns.cpu(|| GpuChain::new(chain.clone(), &gpu.device, width, height));
2695 for (label, rect) in std::iter::once(("full", None))
2696 .chain(rects.iter().map(|(label, rect)| (*label, Some([*rect]))))
2697 {
2698 let mut samples = Vec::new();
2699 for attempt in 0..15 {
2700 let mut encoder =
2701 gpu.device
2702 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
2703 label: Some("region benchmark"),
2704 });
2705 match &rect {
2706 Some(rects) => {
2707 post.run_region(
2708 &mut encoder,
2709 &input.view,
2710 None,
2711 None,
2712 &output.view,
2713 0,
2714 0,
2715 1.0,
2716 rects,
2717 Some(&mut timer),
2718 );
2719 }
2720 None => post.run_timed(
2721 &mut encoder,
2722 &input.view,
2723 None,
2724 None,
2725 &output.view,
2726 0,
2727 0,
2728 Some(&mut timer),
2729 ),
2730 }
2731 let slot = timer.finish(&mut encoder);
2732 gpu.queue.submit(Some(encoder.finish()));
2733 if let Some(slot) = slot {
2734 timer.submitted(slot);
2735 }
2736 gpu.device
2737 .poll(wgpu::PollType::wait_indefinitely())
2738 .unwrap();
2739 let timings = timer.collect(&gpu.device);
2740 let total: f64 = timings.iter().flatten().map(|pass| pass.milliseconds).sum();
2741 turns.add(total);
2742 if attempt > 0 {
2743 samples.push(total);
2744 }
2745 }
2746 samples.sort_by(f64::total_cmp);
2747 println!("{name} {label} {:.4} ms", samples[7]);
2748 }
2749 }
2750 }
2751
2752 pub(crate) const CAPACITY: (u32, u32) = (160, 96);
2753 pub(crate) const VIEWPORTS: [(u32, u32); 2] = [(120, 72), (117, 71)];
2754 const LOUD: [f32; 4] = [400.0, 0.0, 900.0, 1.0];
2755 const BENT: [f32; 4] = [5.0, -7.0, 3.0, 1.0];
2756 const SENTINEL: f32 = 77.0;
2757
2758 #[test]
2759 fn a_viewport_refuses_empty_and_oversized_rects() {
2760 assert_eq!(fit_viewport((160, 96), (160, 96)), Ok((160, 96)));
2761 assert_eq!(fit_viewport((1, 1), (160, 96)), Ok((1, 1)));
2762 assert_eq!(fit_viewport((117, 71), (160, 96)), Ok((117, 71)));
2763 for viewport in [(0, 96), (160, 0), (0, 0), (161, 96), (160, 97), (400, 400)] {
2764 assert!(
2765 fit_viewport(viewport, (160, 96)).is_err(),
2766 "{viewport:?} was accepted"
2767 );
2768 }
2769 }
2770
2771 fn detail(x: u32, y: u32) -> [f32; 4] {
2772 let shades = [0.0, 0.25, 0.5, 0.75, 1.0, 2.0];
2773 let fine = ((x * 7 + y * 3) % 17) as f32 * 0.02;
2774 [
2775 shades[((x / 8 + y / 16) % 6) as usize] + fine,
2776 shades[((x / 16 + y / 8 + 2) % 6) as usize],
2777 shades[((x / 8 + 3) % 6) as usize] + fine * 0.5,
2778 1.0,
2779 ]
2780 }
2781
2782 fn tilted(x: u32, y: u32) -> [f32; 4] {
2783 match (x / 9 + y / 13) % 3 {
2784 0 => [0.0, 0.6, 0.8, 1.0],
2785 1 => [0.7, 0.0, 0.71, 1.0],
2786 _ => [0.0, 0.0, 1.0, 1.0],
2787 }
2788 }
2789
2790 struct Scene {
2791 input: OffscreenTarget,
2792 normal: OffscreenTarget,
2793 _depth: wgpu::Texture,
2794 depth: wgpu::TextureView,
2795 }
2796
2797 pub(crate) fn framed(
2798 size: (u32, u32),
2799 active: (u32, u32),
2800 inside: fn(u32, u32) -> [f32; 4],
2801 outside: [f32; 4],
2802 ) -> Vec<u16> {
2803 Image::from_fn(size.0, size.1, |x, y| {
2804 if x < active.0 && y < active.1 {
2805 inside(x, y)
2806 } else {
2807 outside
2808 }
2809 })
2810 .pixels
2811 .iter()
2812 .flat_map(|p| p.iter().map(|v| half(*v)))
2813 .collect()
2814 }
2815
2816 fn depth_texture(
2817 gpu: &Gpu,
2818 size: (u32, u32),
2819 active: (u32, u32),
2820 ) -> (wgpu::Texture, wgpu::TextureView) {
2821 let texture = gpu.device.create_texture(&wgpu::TextureDescriptor {
2822 label: Some("post viewport depth"),
2823 size: wgpu::Extent3d {
2824 width: size.0,
2825 height: size.1,
2826 depth_or_array_layers: 1,
2827 },
2828 mip_level_count: 1,
2829 sample_count: 1,
2830 dimension: wgpu::TextureDimension::D2,
2831 format: wgpu::TextureFormat::Depth32Float,
2832 usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::RENDER_ATTACHMENT,
2833 view_formats: &[],
2834 });
2835 let view = texture.create_view(&wgpu::TextureViewDescriptor::default());
2836 let source = format!(
2837 "@vertex fn vs(@builtin(vertex_index) i: u32) -> @builtin(position) vec4<f32> {{ return vec4<f32>(f32((i << 1u) & 2u) * 2.0 - 1.0, f32(i & 2u) * 2.0 - 1.0, 0.0, 1.0); }} @fragment fn fs(@builtin(position) at: vec4<f32>) -> @builtin(frag_depth) f32 {{ let p = vec2<u32>(at.xy); if (p.x >= {}u || p.y >= {}u) {{ return 0.02; }} return select(0.35, 0.7, ((p.x / 11u + p.y / 7u) % 2u) == 1u) + f32(p.x % 5u) * 0.01; }}",
2838 active.0, active.1
2839 );
2840 let module = gpu
2841 .device
2842 .create_shader_module(wgpu::ShaderModuleDescriptor {
2843 label: Some("post viewport depth"),
2844 source: wgpu::ShaderSource::Wgsl(source.into()),
2845 });
2846 let pipeline = gpu
2847 .device
2848 .create_render_pipeline(&wgpu::RenderPipelineDescriptor {
2849 label: Some("post viewport depth"),
2850 layout: None,
2851 vertex: wgpu::VertexState {
2852 module: &module,
2853 entry_point: Some("vs"),
2854 compilation_options: wgpu::PipelineCompilationOptions::default(),
2855 buffers: &[],
2856 },
2857 primitive: wgpu::PrimitiveState::default(),
2858 depth_stencil: Some(wgpu::DepthStencilState {
2859 format: wgpu::TextureFormat::Depth32Float,
2860 depth_write_enabled: Some(true),
2861 depth_compare: Some(wgpu::CompareFunction::Always),
2862 stencil: wgpu::StencilState::default(),
2863 bias: wgpu::DepthBiasState::default(),
2864 }),
2865 multisample: wgpu::MultisampleState::default(),
2866 fragment: Some(wgpu::FragmentState {
2867 module: &module,
2868 entry_point: Some("fs"),
2869 compilation_options: wgpu::PipelineCompilationOptions::default(),
2870 targets: &[],
2871 }),
2872 multiview_mask: None,
2873 cache: None,
2874 });
2875 let mut encoder = gpu
2876 .device
2877 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
2878 label: Some("post viewport depth"),
2879 });
2880 {
2881 let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
2882 multiview_mask: None,
2883 label: Some("post viewport depth"),
2884 color_attachments: &[],
2885 depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
2886 view: &view,
2887 depth_ops: Some(wgpu::Operations {
2888 load: wgpu::LoadOp::Clear(1.0),
2889 store: wgpu::StoreOp::Store,
2890 }),
2891 stencil_ops: None,
2892 }),
2893 timestamp_writes: None,
2894 occlusion_query_set: None,
2895 });
2896 pass.set_pipeline(&pipeline);
2897 pass.draw(0..3, 0..1);
2898 }
2899 gpu.queue.submit(Some(encoder.finish()));
2900 (texture, view)
2901 }
2902
2903 fn scene(gpu: &Gpu, size: (u32, u32), active: (u32, u32)) -> Scene {
2904 let input = gpu.offscreen(size.0, size.1, FORMAT).unwrap();
2905 gpu.upload_rgba16(&input, &framed(size, active, detail, LOUD))
2906 .unwrap();
2907 let normal = gpu.offscreen(size.0, size.1, FORMAT).unwrap();
2908 gpu.upload_rgba16(&normal, &framed(size, active, tilted, BENT))
2909 .unwrap();
2910 let (texture, depth) = depth_texture(gpu, size, active);
2911 Scene {
2912 input,
2913 normal,
2914 _depth: texture,
2915 depth,
2916 }
2917 }
2918
2919 fn sentinel(gpu: &Gpu, size: (u32, u32)) -> OffscreenTarget {
2920 let output = output_target(gpu, size.0, size.1);
2921 let fill = vec![half(SENTINEL); (size.0 * size.1 * 4) as usize];
2922 gpu.upload_rgba16(&output, &fill).unwrap();
2923 output
2924 }
2925
2926 fn post_scene(
2927 gpu: &Gpu,
2928 post: &GpuChain,
2929 scene: &Scene,
2930 output: &OffscreenTarget,
2931 chain: &Chain,
2932 rects: Option<&[[u32; 4]]>,
2933 ) -> Vec<u16> {
2934 let mut encoder = gpu
2935 .device
2936 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
2937 label: Some("post viewport test"),
2938 });
2939 match rects {
2940 Some(rects) => {
2941 post.run_region(
2942 &mut encoder,
2943 &scene.input.view,
2944 Some(&scene.depth),
2945 Some(&scene.normal.view),
2946 &output.view,
2947 chain.frame,
2948 chain.seed,
2949 1.0,
2950 rects,
2951 None,
2952 );
2953 }
2954 None => post.run(
2955 &mut encoder,
2956 &scene.input.view,
2957 Some(&scene.depth),
2958 Some(&scene.normal.view),
2959 &output.view,
2960 chain.frame,
2961 chain.seed,
2962 ),
2963 }
2964 gpu.queue.submit(Some(encoder.finish()));
2965 gpu.readback_rgba16(output).unwrap()
2966 }
2967
2968 fn viewport_chains() -> Vec<(&'static str, Chain)> {
2969 let chain = |passes: Vec<Pass>| {
2970 let mut passes = passes;
2971 passes.push(Pass::Encode);
2972 Chain {
2973 passes,
2974 seed: 31,
2975 frame: 5,
2976 }
2977 };
2978 let mut chains = vec![
2979 ("standard finish", as_live(frame_finish())),
2980 ("tent finish", chain(fused_finish(0.5, 0.3))),
2981 (
2982 "linear bloom",
2983 chain(vec![
2984 Pass::Tone(crate::tone::Tone::aces()),
2985 Pass::Bloom(crate::bloom::Bloom::neutral(0.45)),
2986 ]),
2987 ),
2988 (
2989 "rings",
2990 chain(vec![
2991 Pass::Tone(crate::tone::Tone::aces()),
2992 Pass::Bloom(crate::bloom::Bloom::rings(
2993 0.5,
2994 0.6,
2995 4,
2996 crate::bloom::NEUTRAL_RADII,
2997 crate::bloom::NEUTRAL_GAINS,
2998 )),
2999 ]),
3000 ),
3001 (
3002 "aberration",
3003 chain(vec![Pass::Exposure(1.0), Pass::Aberration(2.5)]),
3004 ),
3005 (
3006 "distortion",
3007 chain(vec![Pass::Exposure(1.0), Pass::Distortion(0.08)]),
3008 ),
3009 (
3010 "pixelate",
3011 chain(vec![Pass::Pixel(crate::fx::Pixelate {
3012 size: 5,
3013 levels: 0,
3014 palette: false,
3015 })]),
3016 ),
3017 (
3018 "scratches",
3019 chain(vec![Pass::Scratches(crate::fx::Scratches {
3020 count: 9,
3021 strength: 0.4,
3022 })]),
3023 ),
3024 (
3025 "halation",
3026 chain(vec![
3027 Pass::Halation(crate::fx::Halation {
3028 strength: 0.6,
3029 threshold: 0.5,
3030 radius: 6,
3031 }),
3032 Pass::Tone(crate::tone::Tone::aces()),
3033 ]),
3034 ),
3035 (
3036 "neon",
3037 chain(vec![Pass::Neon(crate::fx::Neon {
3038 saturation: 1.8,
3039 strength: 1.35,
3040 radius: 3,
3041 })]),
3042 ),
3043 ];
3044 chains.extend(region_chains());
3045 chains.extend(
3046 Style::ALL
3047 .into_iter()
3048 .map(|style| (style.name(), as_live(style.chain()))),
3049 );
3050 let mut tape = as_live(Style::Crt.chain());
3051 tape.set_tape(crate::tape::Tape::rewind(1.0));
3052 tape.frame = 30;
3053 tape.seed = 7;
3054 chains.push(("tape", tape));
3055 chains
3056 }
3057
3058 fn samples_in_hardware(post: &GpuChain) -> bool {
3059 post.active()
3060 .iter()
3061 .flat_map(|group| &group.stages)
3062 .any(|stage| {
3063 matches!(
3064 stage.aux,
3065 Aux::BloomLinear | Aux::LinearSampler | Aux::GradeBloomLinear
3066 )
3067 })
3068 }
3069
3070 fn compare_viewport(
3071 name: &str,
3072 got: &[u16],
3073 want: &[u16],
3074 viewport: (u32, u32),
3075 hardware: bool,
3076 ) -> f32 {
3077 let (w, h) = viewport;
3078 let mut worst = 0.0f32;
3079 for y in 0..CAPACITY.1 {
3080 for x in 0..CAPACITY.0 {
3081 let at = ((y * CAPACITY.0 + x) * 4) as usize;
3082 if x >= w || y >= h {
3083 assert!(
3084 got[at..at + 4].iter().all(|v| *v == half(SENTINEL)),
3085 "{name} {viewport:?} wrote ({x},{y}) outside its viewport"
3086 );
3087 continue;
3088 }
3089 let there = ((y * w + x) * 4) as usize;
3090 for channel in 0..4 {
3091 let a = got[at + channel];
3092 let b = want[there + channel];
3093 let delta = if a == b {
3094 0.0
3095 } else {
3096 (float(a) - float(b)).abs()
3097 };
3098 let delta = if delta.is_nan() { f32::INFINITY } else { delta };
3099 worst = worst.max(delta);
3100 if hardware {
3101 assert!(
3102 delta <= 1.0 / 255.0,
3103 "{name} {viewport:?} ({x},{y}) channel {channel} must be within 1/255 of a chain at its size: {} vs {}",
3104 float(a),
3105 float(b)
3106 );
3107 } else {
3108 assert!(
3109 a == b,
3110 "{name} {viewport:?} ({x},{y}) channel {channel} must equal a chain at its size exactly: {} vs {}",
3111 float(a),
3112 float(b)
3113 );
3114 }
3115 }
3116 }
3117 }
3118 worst
3119 }
3120
3121 #[test]
3122 #[ignore = "needs a GPU; run under pgpu"]
3123 fn a_viewport_post_equals_a_chain_at_the_viewport_size() {
3124 let gpu = pollster::block_on(Gpu::headless()).unwrap();
3125 let garbage = scene(&gpu, CAPACITY, (0, 0));
3126 let scenes: Vec<(Scene, Scene)> = VIEWPORTS
3127 .iter()
3128 .map(|&viewport| {
3129 (
3130 scene(&gpu, CAPACITY, viewport),
3131 scene(&gpu, viewport, viewport),
3132 )
3133 })
3134 .collect();
3135 let mut checked = 0;
3136 for (name, chain) in viewport_chains() {
3137 let mut post = GpuChain::new(chain.clone(), &gpu.device, CAPACITY.0, CAPACITY.1);
3138 let hardware = samples_in_hardware(&post);
3139 let mut report = Vec::new();
3140 for (&viewport, (framed, exact)) in VIEWPORTS.iter().zip(&scenes) {
3141 post.set_viewport(CAPACITY.0, CAPACITY.1).unwrap();
3142 post_scene(
3143 &gpu,
3144 &post,
3145 &garbage,
3146 &output_target(&gpu, CAPACITY.0, CAPACITY.1),
3147 &chain,
3148 None,
3149 );
3150 post.set_viewport(viewport.0, viewport.1).unwrap();
3151 assert_eq!(post.viewport(), viewport);
3152 assert_eq!(post.capacity(), CAPACITY);
3153 let got = post_scene(&gpu, &post, framed, &sentinel(&gpu, CAPACITY), &chain, None);
3154 let reference = GpuChain::new(chain.clone(), &gpu.device, viewport.0, viewport.1);
3155 let want = post_scene(
3156 &gpu,
3157 &reference,
3158 exact,
3159 &output_target(&gpu, viewport.0, viewport.1),
3160 &chain,
3161 None,
3162 );
3163 let worst = compare_viewport(name, &got, &want, viewport, hardware);
3164 report.push(format!("{viewport:?} {:.5}", worst * 255.0));
3165 if post.region_reach().is_some() {
3166 let rects = [[0, 0, CAPACITY.0, CAPACITY.1]];
3167 let region = post_scene(
3168 &gpu,
3169 &post,
3170 framed,
3171 &sentinel(&gpu, CAPACITY),
3172 &chain,
3173 Some(&rects),
3174 );
3175 let worst = compare_viewport(name, ®ion, &want, viewport, hardware);
3176 report.push(format!("region {:.5}", worst * 255.0));
3177 }
3178 checked += 1;
3179 }
3180 println!(
3181 "{name} ({}): worst {} /255",
3182 if hardware { "sampled" } else { "exact" },
3183 report.join(", ")
3184 );
3185 }
3186 assert!(checked >= 2 * 30, "{checked}");
3187 }
3188
3189 #[test]
3190 #[ignore = "needs a GPU; run under pgpu"]
3191 fn a_full_viewport_post_is_byte_identical() {
3192 let gpu = pollster::block_on(Gpu::headless()).unwrap();
3193 let garbage = scene(&gpu, CAPACITY, (0, 0));
3194 let whole = scene(&gpu, CAPACITY, CAPACITY);
3195 for (name, chain) in viewport_chains() {
3196 let fresh = GpuChain::new(chain.clone(), &gpu.device, CAPACITY.0, CAPACITY.1);
3197 let want = post_scene(
3198 &gpu,
3199 &fresh,
3200 &whole,
3201 &output_target(&gpu, CAPACITY.0, CAPACITY.1),
3202 &chain,
3203 None,
3204 );
3205 let mut post = GpuChain::new(chain.clone(), &gpu.device, CAPACITY.0, CAPACITY.1);
3206 post.set_viewport(VIEWPORTS[1].0, VIEWPORTS[1].1).unwrap();
3207 post_scene(
3208 &gpu,
3209 &post,
3210 &garbage,
3211 &output_target(&gpu, CAPACITY.0, CAPACITY.1),
3212 &chain,
3213 None,
3214 );
3215 post.set_viewport(CAPACITY.0, CAPACITY.1).unwrap();
3216 let got = post_scene(
3217 &gpu,
3218 &post,
3219 &whole,
3220 &output_target(&gpu, CAPACITY.0, CAPACITY.1),
3221 &chain,
3222 None,
3223 );
3224 assert!(got == want, "{name} changed bytes at the full viewport");
3225 post.set_viewport(VIEWPORTS[0].0, VIEWPORTS[0].1).unwrap();
3226 post.resize(CAPACITY.0, CAPACITY.1);
3227 assert_eq!(post.viewport(), CAPACITY, "{name}");
3228 }
3229 }
3230}