1use std::collections::HashMap;
2use std::path::PathBuf;
3use std::sync::atomic::{AtomicU64, Ordering};
4use std::sync::mpsc;
5use std::time::Instant;
6
7use crate::wgpu;
8
9mod frames;
10pub mod trace;
11
12pub use frames::{
13 CAP_FPS, FrameClock, FrameGpu, Frames, Pace, Sink, TURN_AFTER_MS, TURN_MS_VAR, TURN_VAR,
14 TURN_WALL_VAR, Tally, TurnReport, UNCAPPED_VAR, WallClock, frozen_file_ms, sleep_to,
15 turn_wall_ms, uncapped,
16};
17pub use trace::TRACE_VAR;
18use trace::{Source, Sources};
19
20#[derive(Clone, Copy, Debug, PartialEq, Eq)]
21pub enum Work {
22 Bands {
23 width: u32,
24 height: u32,
25 },
26 Tiles {
27 width: u32,
28 height: u32,
29 tile_width: u32,
30 },
31 Dispatches {
32 count: u32,
33 },
34}
35
36#[derive(Clone, Copy, Debug, PartialEq, Eq)]
37pub enum Slice {
38 Pixels {
39 x: u32,
40 y: u32,
41 width: u32,
42 height: u32,
43 },
44 Dispatches {
45 start: u32,
46 count: u32,
47 },
48}
49
50#[derive(Clone, Copy, Debug, PartialEq, Eq)]
51pub enum Timing {
52 Gpu,
53 Cpu,
54}
55
56#[derive(Clone, Copy, Debug, Default, PartialEq, Eq)]
57pub enum TimingSource {
58 #[default]
59 Unmeasured,
60 Gpu,
61 Cpu,
62 Mixed,
63}
64
65#[derive(Clone, Debug, Default)]
66pub struct Stats {
67 pub milliseconds: Vec<f64>,
68 pub timings: Vec<Timing>,
69 pub wall_ms: f64,
70}
71
72impl Stats {
73 pub fn count(&self) -> usize {
74 self.milliseconds.len()
75 }
76
77 pub fn longest_ms(&self) -> f64 {
78 self.milliseconds.iter().copied().fold(0.0, f64::max)
79 }
80
81 pub fn median_ms(&self) -> f64 {
82 let mut values = self.milliseconds.clone();
83 values.sort_by(f64::total_cmp);
84 if values.is_empty() {
85 return 0.0;
86 }
87 let middle = values.len() / 2;
88 if values.len().is_multiple_of(2) {
89 (values[middle - 1] + values[middle]) / 2.0
90 } else {
91 values[middle]
92 }
93 }
94
95 pub fn percentile_ms(&self, percent: f64) -> f64 {
96 let mut values = self.milliseconds.clone();
97 values.sort_by(f64::total_cmp);
98 if values.is_empty() {
99 return 0.0;
100 }
101 let rank = (percent.clamp(0.0, 100.0) / 100.0 * values.len() as f64).ceil() as usize;
102 values[rank.clamp(1, values.len()) - 1]
103 }
104
105 pub fn source(&self) -> TimingSource {
106 let gpu = self.timings.contains(&Timing::Gpu);
107 let cpu = self.timings.contains(&Timing::Cpu);
108 match (gpu, cpu) {
109 (false, false) => TimingSource::Unmeasured,
110 (true, false) => TimingSource::Gpu,
111 (false, true) => TimingSource::Cpu,
112 (true, true) => TimingSource::Mixed,
113 }
114 }
115}
116
117#[derive(Clone, Copy, Debug, Default, PartialEq)]
118pub struct Pacing {
119 pub work_ms: f64,
120 pub longest_ms: f64,
121 pub slices: usize,
122 pub untimed: usize,
123 pub source: TimingSource,
124}
125
126impl Pacing {
127 fn add(&mut self, milliseconds: f64, timing: Timing, reported: bool) {
128 if !reported {
129 self.untimed += 1;
130 } else if milliseconds.is_finite() && milliseconds >= 0.0 {
131 self.work_ms += milliseconds;
132 self.longest_ms = self.longest_ms.max(milliseconds);
133 }
134 self.slices += 1;
135 self.source = match (self.source, timing) {
136 (TimingSource::Unmeasured, Timing::Gpu) | (TimingSource::Gpu, Timing::Gpu) => {
137 TimingSource::Gpu
138 }
139 (TimingSource::Unmeasured, Timing::Cpu) | (TimingSource::Cpu, Timing::Cpu) => {
140 TimingSource::Cpu
141 }
142 _ => TimingSource::Mixed,
143 };
144 }
145}
146
147pub trait Turned {
148 fn turned(self) -> bool;
149}
150
151impl Turned for () {
152 fn turned(self) -> bool {
153 false
154 }
155}
156
157impl Turned for bool {
158 fn turned(self) -> bool {
159 self
160 }
161}
162
163pub fn turn_args(work_ms: f64, longest_ms: f64) -> [String; 5] {
164 [
165 "turn".into(),
166 "--work-ms".into(),
167 format!("{work_ms:.3}"),
168 "--longest-ms".into(),
169 format!("{longest_ms:.3}"),
170 ]
171}
172
173pub const FROZEN_VAR: &str = "GPU_QUEUE_FROZEN";
174
175#[derive(Clone, Debug, Default, PartialEq, Eq)]
176pub struct Frozen {
177 path: Option<PathBuf>,
178}
179
180impl Frozen {
181 pub fn from_env() -> Self {
182 Self {
183 path: std::env::var_os(FROZEN_VAR)
184 .filter(|path| !path.is_empty())
185 .map(PathBuf::from),
186 }
187 }
188
189 pub fn at(path: impl Into<PathBuf>) -> Self {
190 Self {
191 path: Some(path.into()),
192 }
193 }
194
195 pub fn none() -> Self {
196 Self::default()
197 }
198
199 pub fn ms(&self) -> f64 {
200 self.path
201 .as_ref()
202 .and_then(|path| std::fs::read_to_string(path).ok())
203 .and_then(|text| text.trim().parse::<f64>().ok())
204 .filter(|ms| ms.is_finite() && *ms >= 0.0)
205 .unwrap_or(0.0)
206 }
207}
208
209pub fn unfrozen_ms(wall_ms: f64, frozen_before_ms: f64, frozen_after_ms: f64) -> f64 {
210 let frozen = (frozen_after_ms - frozen_before_ms).max(0.0);
211 (wall_ms - frozen).max(0.0)
212}
213
214#[derive(Clone, Copy, Debug, PartialEq)]
215pub struct Moment {
216 pub at: Instant,
217 pub frozen_ms: f64,
218}
219
220impl Moment {
221 pub fn now(frozen: &Frozen) -> Self {
222 Self {
223 at: Instant::now(),
224 frozen_ms: frozen.ms(),
225 }
226 }
227
228 pub fn ms_since(&self, earlier: Moment) -> f64 {
229 unfrozen_ms(
230 self.at.saturating_duration_since(earlier.at).as_secs_f64() * 1000.0,
231 earlier.frozen_ms,
232 self.frozen_ms,
233 )
234 }
235
236 fn later(self, other: Moment) -> Moment {
237 if other.at > self.at { other } else { self }
238 }
239}
240
241pub fn gpu_ms(passes: &[crate::PassTiming]) -> Option<f64> {
242 let total: f64 = passes.iter().map(|pass| pass.milliseconds).sum();
243 (!passes.is_empty() && total.is_finite() && total > 0.0).then_some(total)
244}
245
246impl crate::FrameTimings {
247 pub fn gpu_ms(&self) -> Option<f64> {
248 gpu_ms(&self.passes)
249 }
250}
251
252static TURNING: std::sync::Mutex<()> = std::sync::Mutex::new(());
253static TURNS_TAKEN: AtomicU64 = AtomicU64::new(0);
254static LAST_TURN: std::sync::Mutex<Option<Instant>> = std::sync::Mutex::new(None);
255
256fn last_turn_or(now: Instant) -> Instant {
257 LAST_TURN
258 .lock()
259 .unwrap_or_else(|poisoned| poisoned.into_inner())
260 .unwrap_or(now)
261}
262
263pub fn turns_taken() -> u64 {
264 TURNS_TAKEN.load(Ordering::SeqCst)
265}
266
267pub const TURN_WALL_MS: f64 = 400.0;
268
269fn env_ms(name: &str, default: f64) -> f64 {
270 std::env::var(name)
271 .ok()
272 .and_then(|value| value.parse::<f64>().ok())
273 .filter(|value| value.is_finite() && *value > 0.0)
274 .unwrap_or(default)
275}
276
277pub struct Turns {
278 queued: bool,
279 threshold_ms: f64,
280 wall_ms: f64,
281 work_ms: f64,
282 longest_ms: f64,
283 measured: bool,
284 untimed: bool,
285 last_turn: Instant,
286 last_event: Instant,
287 interval_ms: f64,
288 taken: u32,
289 frozen: Frozen,
290 site: &'static std::panic::Location<'static>,
291 sources: Sources,
292}
293
294impl Default for Turns {
295 #[track_caller]
296 fn default() -> Self {
297 let queued =
298 std::env::var_os(TURN_VAR).is_some() || std::env::var_os("GPU_QUEUE_PID").is_some();
299 Self::new(env_ms(TURN_MS_VAR, 50.0), queued)
300 .with_wall_ms(env_ms(TURN_WALL_VAR, TURN_WALL_MS))
301 }
302}
303
304impl Turns {
305 #[track_caller]
306 pub fn new(threshold_ms: f64, queued: bool) -> Self {
307 trace::arm();
308 trace::created(std::panic::Location::caller());
309 Self {
310 queued,
311 threshold_ms: if threshold_ms.is_finite() {
312 threshold_ms.max(1.0)
313 } else {
314 50.0
315 },
316 wall_ms: TURN_WALL_MS,
317 work_ms: 0.0,
318 longest_ms: 0.0,
319 measured: false,
320 untimed: false,
321 last_turn: last_turn_or(Instant::now()),
322 last_event: Instant::now(),
323 interval_ms: 0.0,
324 taken: 0,
325 frozen: Frozen::from_env(),
326 site: std::panic::Location::caller(),
327 sources: Sources::default(),
328 }
329 }
330
331 pub fn with_frozen(mut self, frozen: Frozen) -> Self {
332 self.frozen = frozen;
333 self
334 }
335
336 pub fn with_wall_ms(mut self, wall_ms: f64) -> Self {
337 if wall_ms.is_finite() && wall_ms > 0.0 {
338 self.wall_ms = wall_ms;
339 }
340 self
341 }
342
343 #[track_caller]
344 pub fn add(&mut self, submission_ms: f64) -> bool {
345 self.site = std::panic::Location::caller();
346 let paced = trace::take_slice();
347 let source = match paced {
348 _ if !(submission_ms.is_finite() && submission_ms >= 0.0) => Source::Untimed,
349 Some(Timing::Gpu) => Source::Gpu,
350 Some(Timing::Cpu) => Source::Wall,
351 None => Source::Given,
352 };
353 self.add_sourced(submission_ms, Instant::now(), source)
354 }
355
356 #[cfg(test)]
357 fn add_at(&mut self, submission_ms: f64, now: Instant) -> bool {
358 self.add_sourced(submission_ms, now, Source::Given)
359 }
360
361 fn noted(&mut self, now: Instant) {
362 self.interval_ms = (now.saturating_duration_since(self.last_event).as_secs_f64() * 1000.0)
363 .min(self.wall_ms / 2.0);
364 self.last_event = now;
365 }
366
367 fn add_sourced(&mut self, submission_ms: f64, now: Instant, source: Source) -> bool {
368 self.noted(now);
369 self.sources.add(source);
370 if submission_ms.is_finite() && submission_ms >= 0.0 {
371 self.work_ms += submission_ms;
372 self.longest_ms = self.longest_ms.max(submission_ms);
373 self.measured = true;
374 } else {
375 self.untimed = true;
376 }
377 if !self.due_at(now) {
378 return false;
379 }
380 self.turn_at(now);
381 true
382 }
383
384 fn due_at(&self, now: Instant) -> bool {
385 self.work_ms >= self.threshold_ms
386 || now.saturating_duration_since(self.last_turn).as_secs_f64() * 1000.0
387 + self.interval_ms
388 >= self.wall_ms
389 }
390
391 #[track_caller]
392 pub fn time<R>(&mut self, submission: impl FnOnce() -> R) -> R {
393 self.site = std::panic::Location::caller();
394 let began = Moment::now(&self.frozen);
395 let turned = turns_taken();
396 let result = submission();
397 let ended = Moment::now(&self.frozen);
398 self.add_timed(began, ended, turns_taken() != turned);
399 result
400 }
401
402 fn add_timed(&mut self, began: Moment, ended: Moment, turned_inside: bool) -> bool {
403 let milliseconds = if turned_inside {
404 f64::NAN
405 } else {
406 ended.ms_since(began)
407 };
408 let source = if turned_inside {
409 Source::Untimed
410 } else {
411 Source::Wall
412 };
413 self.add_sourced(milliseconds, ended.at, source)
414 }
415
416 #[track_caller]
417 pub fn cpu<R>(&mut self, work: impl FnOnce() -> R) -> R {
418 self.poll();
419 let result = work();
420 self.poll();
421 result
422 }
423
424 #[track_caller]
425 pub fn poll(&mut self) -> bool {
426 self.site = std::panic::Location::caller();
427 let now = Instant::now();
428 self.noted(now);
429 if !self.due_at(now) {
430 return false;
431 }
432 self.turn_at(now);
433 true
434 }
435
436 #[track_caller]
437 pub fn turn(&mut self) {
438 self.site = std::panic::Location::caller();
439 self.turn_at(Instant::now());
440 }
441
442 pub fn args(&self) -> Vec<String> {
443 if self.untimed && !self.measured {
444 vec!["turn".into()]
445 } else {
446 turn_args(self.work_ms, self.longest_ms).to_vec()
447 }
448 }
449
450 fn turn_at(&mut self, now: Instant) {
451 let args = self.args();
452 trace::record(
453 "turns",
454 self.site,
455 self.work_ms,
456 self.longest_ms,
457 self.sources,
458 );
459 self.sources = Sources::default();
460 self.work_ms = 0.0;
461 self.longest_ms = 0.0;
462 self.measured = false;
463 self.untimed = false;
464 self.last_turn = now;
465 self.last_event = now;
466 self.taken += 1;
467 TURNS_TAKEN.fetch_add(1, Ordering::SeqCst);
468 if self.queued {
469 let _held = TURNING
470 .lock()
471 .unwrap_or_else(|poisoned| poisoned.into_inner());
472 let _ = std::process::Command::new("pgpu")
473 .args(args)
474 .stdout(std::process::Stdio::null())
475 .status();
476 }
477 trace::turned();
478 *LAST_TURN
479 .lock()
480 .unwrap_or_else(|poisoned| poisoned.into_inner()) = Some(Instant::now());
481 }
482
483 pub fn pending(&self) -> (f64, f64) {
484 (self.work_ms, self.longest_ms)
485 }
486
487 pub fn taken(&self) -> u32 {
488 self.taken
489 }
490}
491
492impl Drop for Turns {
493 fn drop(&mut self) {
494 trace::record(
495 "end",
496 self.site,
497 self.work_ms,
498 self.longest_ms,
499 self.sources,
500 );
501 }
502}
503
504const DEFAULT_CEILING_MS: f64 = 6.0;
505const CUT_SAFETY: f64 = 0.8;
506const RECOVERY_SLICES: u32 = 6;
507const RECOVERY_GROWTH: f64 = 1.25;
508const DEAD_BAND: f64 = 0.15;
509const WINDOW: usize = 8;
510const DISTINCT: f64 = 0.2;
511const REGIME: f64 = 1.5;
512const OVERHEAD_FLOOR: f64 = 0.05;
513const PER_UNIT_FLOOR: f64 = 0.02;
514const UNKNOWN_OVERHEAD: f64 = 0.3;
515
516fn median(values: &mut [f64]) -> f64 {
517 values.sort_by(f64::total_cmp);
518 let middle = values.len() / 2;
519 if values.len().is_multiple_of(2) {
520 (values[middle - 1] + values[middle]) / 2.0
521 } else {
522 values[middle]
523 }
524}
525
526#[derive(Clone, Copy, Debug)]
527struct Affine {
528 overhead_ms: f64,
529 per_unit_ms: f64,
530}
531
532impl Affine {
533 fn unknown_overhead(weight: f64, elapsed_ms: f64) -> Self {
534 Self {
535 overhead_ms: elapsed_ms * UNKNOWN_OVERHEAD,
536 per_unit_ms: elapsed_ms * (1.0 - UNKNOWN_OVERHEAD) / weight,
537 }
538 }
539
540 fn at(&self, units: f64) -> f64 {
541 self.overhead_ms + self.per_unit_ms * units
542 }
543
544 fn units_for(&self, target_ms: f64) -> f64 {
545 ((target_ms - self.overhead_ms) / self.per_unit_ms).max(1.0)
546 }
547}
548
549#[derive(Clone, Copy, Debug)]
550struct History {
551 points: [(f64, f64); WINDOW],
552 len: usize,
553 at: usize,
554}
555
556impl History {
557 fn new() -> Self {
558 Self {
559 points: [(0.0, 0.0); WINDOW],
560 len: 0,
561 at: 0,
562 }
563 }
564
565 fn points(&self) -> &[(f64, f64)] {
566 &self.points[..self.len]
567 }
568
569 fn push(&mut self, weight: f64, elapsed_ms: f64) {
570 self.points[self.at] = (weight, elapsed_ms);
571 self.at = (self.at + 1) % WINDOW;
572 self.len = (self.len + 1).min(WINDOW);
573 }
574
575 fn restart(&mut self, weight: f64, elapsed_ms: f64) {
576 *self = Self::new();
577 self.push(weight, elapsed_ms);
578 }
579
580 fn fit(&self) -> Option<Affine> {
581 let points = self.points();
582 let mut slopes = Vec::with_capacity(WINDOW * (WINDOW - 1) / 2);
583 for (index, &(first_weight, first_ms)) in points.iter().enumerate() {
584 for &(second_weight, second_ms) in &points[index + 1..] {
585 let widest = first_weight.max(second_weight);
586 if (first_weight - second_weight).abs() >= DISTINCT * widest {
587 slopes.push((second_ms - first_ms) / (second_weight - first_weight));
588 }
589 }
590 }
591 if slopes.is_empty() {
592 return None;
593 }
594 let steepness = median(&mut slopes);
595 let mut weights: Vec<f64> = points.iter().map(|point| point.0).collect();
596 let mut times: Vec<f64> = points.iter().map(|point| point.1).collect();
597 let fastest = times.iter().copied().fold(f64::MAX, f64::min);
598 let floor = PER_UNIT_FLOOR * median(&mut times) / median(&mut weights);
599 let per_unit_ms = steepness.max(floor);
600 let mut offsets: Vec<f64> = points
601 .iter()
602 .map(|&(weight, ms)| ms - per_unit_ms * weight)
603 .collect();
604 Some(Affine {
605 overhead_ms: median(&mut offsets).clamp(0.0, fastest),
606 per_unit_ms,
607 })
608 }
609
610 fn expected_ms(&self, weight: f64) -> Option<f64> {
611 if let Some(model) = self.fit() {
612 return Some(model.at(weight));
613 }
614 let mut weights: Vec<f64> = self.points().iter().map(|point| point.0).collect();
615 if weights.is_empty() {
616 return None;
617 }
618 let usual = median(&mut weights);
619 if (usual - weight).abs() >= DISTINCT * usual.max(weight) {
620 return None;
621 }
622 let mut times: Vec<f64> = self.points().iter().map(|point| point.1).collect();
623 Some(median(&mut times))
624 }
625
626 fn record(&mut self, weight: f64, elapsed_ms: f64, over_ceiling: bool) {
627 let shifted = self.expected_ms(weight).is_some_and(|expected| {
628 elapsed_ms > expected * REGIME || elapsed_ms * REGIME < expected
629 });
630 if over_ceiling || shifted {
631 self.restart(weight, elapsed_ms);
632 } else {
633 self.push(weight, elapsed_ms);
634 }
635 }
636}
637
638#[derive(Clone, Copy, Debug)]
639struct Tuning {
640 units: u32,
641 recovering: u32,
642 history: History,
643}
644
645#[derive(Clone, Copy, Debug)]
646struct Limits {
647 target_ms: f64,
648 ceiling_ms: f64,
649 step: u32,
650}
651
652fn on_the_grid(next: u32, used: u32, step: u32, nearest: bool) -> u32 {
653 if step <= 1 || next == used {
654 return next;
655 }
656 if next < used || !nearest {
657 (next / step * step).max(step)
658 } else {
659 ((next + step / 2) / step * step).max(step)
660 }
661}
662
663impl Tuning {
664 fn new(units: u32) -> Self {
665 Self {
666 units,
667 recovering: 0,
668 history: History::new(),
669 }
670 }
671
672 fn learn(&mut self, used: u32, weight: f64, elapsed_ms: f64, available: u32, limits: Limits) {
673 let available = available.max(1);
674 if !elapsed_ms.is_finite() || elapsed_ms <= 0.0 || weight <= 0.0 {
675 self.units = self.units.clamp(1, available);
676 return;
677 }
678 let over_ceiling = elapsed_ms > limits.ceiling_ms;
679 self.history.record(weight, elapsed_ms, over_ceiling);
680 let fit = self.history.fit();
681 let ideal = match fit {
682 Some(model) if model.overhead_ms >= OVERHEAD_FLOOR * limits.target_ms => {
683 model.units_for(limits.target_ms)
684 }
685 _ => weight * limits.target_ms / elapsed_ms,
686 };
687 let next = if over_ceiling {
688 self.recovering = RECOVERY_SLICES;
689 let next = ((ideal * CUT_SAFETY).floor() as u32)
690 .max(1)
691 .min(self.units)
692 .min(used);
693 on_the_grid(next, used, limits.step, false)
694 } else {
695 let allow_growth = weight * 4.0 >= used as f64;
696 let mut next = adapt(used, ideal, allow_growth);
697 let mut nearest = true;
698 if self.recovering > 0 {
699 self.recovering -= 1;
700 nearest = false;
701 let cap = (self.units as f64 * RECOVERY_GROWTH).ceil() as u32;
702 next = next.min(cap.max(self.units));
703 }
704 let next = on_the_grid(next, used, limits.step, nearest);
705 let model = fit.unwrap_or_else(|| Affine::unknown_overhead(weight, elapsed_ms));
706 let further = used.saturating_add(limits.step);
707 let coarse = limits.step as f64 >= DEAD_BAND * used as f64;
708 if limits.step > 1
709 && next <= used
710 && allow_growth
711 && coarse
712 && ideal > used as f64
713 && model.at(further as f64) <= limits.ceiling_ms
714 {
715 further
716 } else {
717 next
718 }
719 };
720 self.units = next.clamp(1, available);
721 }
722}
723
724pub struct Pacer {
725 target_ms: f64,
726 ceiling_ms: f64,
727 learned: HashMap<String, Tuning>,
728 fixed_units: Option<u32>,
729 max_units: Option<u32>,
730 step: (u32, u32),
731 pacing: Pacing,
732}
733
734impl Default for Pacer {
735 fn default() -> Self {
736 Self::new(3.5)
737 }
738}
739
740impl Pacer {
741 pub fn new(target_ms: f64) -> Self {
742 let target_ms = if target_ms.is_finite() {
743 target_ms.clamp(2.0, 6.0)
744 } else {
745 3.5
746 };
747 Self {
748 target_ms,
749 ceiling_ms: DEFAULT_CEILING_MS.max(target_ms),
750 learned: HashMap::new(),
751 fixed_units: None,
752 max_units: None,
753 step: (1, 1),
754 pacing: Pacing::default(),
755 }
756 }
757
758 pub fn set_step(&mut self, rows: u32, columns: u32) {
759 self.step = (rows.max(1), columns.max(1));
760 }
761
762 pub fn step(&self) -> (u32, u32) {
763 self.step
764 }
765
766 pub fn set_fixed_units(&mut self, units: Option<u32>) {
767 self.fixed_units = units.map(|value| value.max(1));
768 }
769
770 pub fn fixed_units(&self) -> Option<u32> {
771 self.fixed_units
772 }
773
774 pub fn set_max_units(&mut self, units: Option<u32>) {
775 self.max_units = units.map(|value| value.max(1));
776 }
777
778 pub fn pacing(&self) -> Pacing {
779 self.pacing
780 }
781
782 pub fn take_pacing(&mut self) -> Pacing {
783 std::mem::take(&mut self.pacing)
784 }
785
786 pub fn set_target_ms(&mut self, target_ms: f64) {
787 if target_ms.is_finite() {
788 self.target_ms = target_ms.clamp(2.0, 100.0);
789 self.ceiling_ms = self.ceiling_ms.max(self.target_ms);
790 }
791 }
792
793 pub fn set_ceiling_ms(&mut self, ceiling_ms: f64) {
794 if ceiling_ms.is_finite() {
795 self.ceiling_ms = ceiling_ms.clamp(self.target_ms, 100.0);
796 }
797 }
798
799 pub fn ceiling_ms(&self) -> f64 {
800 self.ceiling_ms
801 }
802
803 pub fn learned_units(&self, label: &str) -> Option<u32> {
804 self.learned.get(label).map(|tuning| tuning.units)
805 }
806
807 pub fn target_ms(&self) -> f64 {
808 self.target_ms
809 }
810
811 pub fn teach(&mut self, label: &str, units: u32) {
812 self.learned
813 .insert(label.to_owned(), Tuning::new(units.max(1)));
814 }
815
816 pub fn run<T: Turned>(
817 &mut self,
818 label: &str,
819 device: &wgpu::Device,
820 queue: &wgpu::Queue,
821 work: Work,
822 encode: impl FnMut(&mut wgpu::CommandEncoder, Slice),
823 between: impl FnMut(f64) -> T,
824 ) -> Result<Stats, String> {
825 let timestamps = device
826 .features()
827 .contains(
828 wgpu::Features::TIMESTAMP_QUERY | wgpu::Features::TIMESTAMP_QUERY_INSIDE_ENCODERS,
829 )
830 .then(|| Timestamp::new(device, queue));
831 let mut lane = GpuLane {
832 label,
833 device,
834 queue,
835 encode,
836 frozen: if timestamps.is_some() {
837 Frozen::none()
838 } else {
839 Frozen::from_env()
840 },
841 timestamps,
842 last_done: None,
843 };
844 self.drive(label, &mut lane, work, between)
845 }
846
847 pub fn run_blocking<T: Turned>(
848 &mut self,
849 label: &str,
850 work: Work,
851 mut submit: impl FnMut(Slice) -> Result<(), String>,
852 between: impl FnMut(f64) -> T,
853 ) -> Result<Stats, String> {
854 self.run_measured(label, work, |slice| submit(slice).map(|_| None), between)
855 }
856
857 pub fn run_measured<T: Turned>(
858 &mut self,
859 label: &str,
860 work: Work,
861 submit: impl FnMut(Slice) -> Result<Option<f64>, String>,
862 between: impl FnMut(f64) -> T,
863 ) -> Result<Stats, String> {
864 let mut lane = MeasuredLane {
865 submit,
866 frozen: Frozen::from_env(),
867 };
868 self.drive(label, &mut lane, work, between)
869 }
870
871 fn drive<L: Lane, T: Turned>(
872 &mut self,
873 label: &str,
874 lane: &mut L,
875 work: Work,
876 mut between: impl FnMut(f64) -> T,
877 ) -> Result<Stats, String> {
878 let mut cursor = Cursor::stepped(work, self.step.0, self.step.1)?;
879 let available = cursor.available();
880 let limits = Limits {
881 target_ms: self.target_ms,
882 ceiling_ms: self.ceiling_ms,
883 step: cursor.step_rows,
884 };
885 let mut tuning = self
886 .learned
887 .get(label)
888 .copied()
889 .unwrap_or_else(|| Tuning::new(1));
890 let fixed = self.fixed_units;
891 let mut stats = Stats::default();
892 let start = Instant::now();
893 let mut pending: Option<Pending<L::Ticket>> = None;
894 let mut slot = 0;
895 loop {
896 let units = cursor.grant(
897 fixed
898 .unwrap_or(tuning.units)
899 .min(self.max_units.unwrap_or(u32::MAX)),
900 );
901 let Some(slice) = cursor.next(units) else {
902 break;
903 };
904 let current = Pending {
905 ticket: lane.submit(slot, slice),
906 slot,
907 weight: cursor.weight(slice),
908 units,
909 };
910 if let Some(previous) = pending.replace(current) {
911 let reported = self.settle(
912 label,
913 lane,
914 &previous,
915 &mut tuning,
916 available,
917 limits,
918 &mut stats,
919 )?;
920 if between(reported).turned() {
921 self.pacing = Pacing::default();
922 }
923 trace::take_slice();
924 lane.resumed();
925 }
926 slot ^= 1;
927 }
928 if let Some(last) = pending {
929 let reported = self.settle(
930 label,
931 lane,
932 &last,
933 &mut tuning,
934 available,
935 limits,
936 &mut stats,
937 )?;
938 if between(reported).turned() {
939 self.pacing = Pacing::default();
940 }
941 trace::take_slice();
942 lane.resumed();
943 }
944 stats.wall_ms = start.elapsed().as_secs_f64() * 1000.0;
945 Ok(stats)
946 }
947
948 #[allow(clippy::too_many_arguments)]
949 fn settle<L: Lane>(
950 &mut self,
951 label: &str,
952 lane: &mut L,
953 done: &Pending<L::Ticket>,
954 tuning: &mut Tuning,
955 available: u32,
956 limits: Limits,
957 stats: &mut Stats,
958 ) -> Result<f64, String> {
959 let measured = lane.finish(done.slot, &done.ticket)?;
960 stats.milliseconds.push(measured.milliseconds);
961 stats.timings.push(measured.timing);
962 self.pacing
963 .add(measured.milliseconds, measured.timing, measured.reported);
964 trace::slice_timed(measured.timing);
965 if self.fixed_units.is_none() {
966 tuning.learn(
967 done.units,
968 done.weight,
969 measured.milliseconds,
970 available,
971 limits,
972 );
973 self.learned.insert(label.to_owned(), *tuning);
974 }
975 Ok(if measured.reported {
976 measured.milliseconds
977 } else {
978 f64::NAN
979 })
980 }
981}
982
983struct Measured {
984 milliseconds: f64,
985 timing: Timing,
986 reported: bool,
987}
988
989trait Lane {
990 type Ticket;
991
992 fn submit(&mut self, slot: usize, slice: Slice) -> Self::Ticket;
993
994 fn finish(&mut self, slot: usize, ticket: &Self::Ticket) -> Result<Measured, String>;
995
996 fn resumed(&mut self) {}
997}
998
999struct Pending<T> {
1000 ticket: T,
1001 slot: usize,
1002 weight: f64,
1003 units: u32,
1004}
1005
1006struct GpuLane<'a, F> {
1007 label: &'a str,
1008 device: &'a wgpu::Device,
1009 queue: &'a wgpu::Queue,
1010 timestamps: Option<Timestamp>,
1011 encode: F,
1012 frozen: Frozen,
1013 last_done: Option<Moment>,
1014}
1015
1016fn fence_ms(submitted: Moment, last_done: Option<Moment>, done: Moment) -> f64 {
1017 let began = last_done.map_or(submitted, |last| last.later(submitted));
1018 done.ms_since(began)
1019}
1020
1021impl<F: FnMut(&mut wgpu::CommandEncoder, Slice)> Lane for GpuLane<'_, F> {
1022 type Ticket = (wgpu::SubmissionIndex, Moment);
1023
1024 fn submit(&mut self, slot: usize, slice: Slice) -> Self::Ticket {
1025 let mut encoder = self
1026 .device
1027 .create_command_encoder(&wgpu::CommandEncoderDescriptor {
1028 label: Some(self.label),
1029 });
1030 if let Some(timestamp) = &self.timestamps {
1031 encoder.write_timestamp(×tamp.slots[slot].queries, 0);
1032 }
1033 (self.encode)(&mut encoder, slice);
1034 if let Some(timestamp) = &self.timestamps {
1035 let timing = ×tamp.slots[slot];
1036 encoder.write_timestamp(&timing.queries, 1);
1037 encoder.resolve_query_set(&timing.queries, 0..2, &timing.resolve, 0);
1038 encoder.copy_buffer_to_buffer(&timing.resolve, 0, &timing.read, 0, 16);
1039 }
1040 let submitted = Moment::now(&self.frozen);
1041 (self.queue.submit(Some(encoder.finish())), submitted)
1042 }
1043
1044 fn finish(&mut self, slot: usize, ticket: &Self::Ticket) -> Result<Measured, String> {
1045 let (index, submitted) = ticket;
1046 self.device
1047 .poll(wgpu::PollType::Wait {
1048 submission_index: Some(index.clone()),
1049 timeout: None,
1050 })
1051 .map_err(|error| format!("{error:?}"))?;
1052 let done = Moment::now(&self.frozen);
1053 let first = self.last_done.is_none();
1054 let cpu_ms = fence_ms(*submitted, self.last_done, done);
1055 self.last_done = Some(done);
1056 let gpu_ms = self
1057 .timestamps
1058 .as_ref()
1059 .and_then(|timestamp| timestamp.read_ms(self.device, slot, index.clone()));
1060 Ok(match gpu_ms {
1061 Some(milliseconds) => Measured {
1062 milliseconds,
1063 timing: Timing::Gpu,
1064 reported: true,
1065 },
1066 None => Measured {
1067 milliseconds: cpu_ms,
1068 timing: Timing::Cpu,
1069 reported: !first,
1070 },
1071 })
1072 }
1073
1074 fn resumed(&mut self) {
1075 if self.last_done.is_some() {
1076 self.last_done = Some(Moment::now(&self.frozen));
1077 }
1078 }
1079}
1080
1081struct MeasuredLane<F> {
1082 submit: F,
1083 frozen: Frozen,
1084}
1085
1086impl<F: FnMut(Slice) -> Result<Option<f64>, String>> Lane for MeasuredLane<F> {
1087 type Ticket = Result<(Option<f64>, f64), String>;
1088
1089 fn submit(&mut self, _slot: usize, slice: Slice) -> Self::Ticket {
1090 let began = Moment::now(&self.frozen);
1091 let gpu_ms = (self.submit)(slice)?;
1092 let cpu_ms = Moment::now(&self.frozen).ms_since(began);
1093 Ok((gpu_ms.filter(|ms| ms.is_finite() && *ms >= 0.0), cpu_ms))
1094 }
1095
1096 fn finish(&mut self, _slot: usize, ticket: &Self::Ticket) -> Result<Measured, String> {
1097 let (gpu_ms, cpu_ms) = ticket.clone()?;
1098 Ok(match gpu_ms {
1099 Some(milliseconds) => Measured {
1100 milliseconds,
1101 timing: Timing::Gpu,
1102 reported: true,
1103 },
1104 None => Measured {
1105 milliseconds: cpu_ms,
1106 timing: Timing::Cpu,
1107 reported: false,
1108 },
1109 })
1110 }
1111}
1112
1113fn adapt(units: u32, ideal: f64, allow_growth: bool) -> u32 {
1114 let ratio = (ideal / units as f64).clamp(0.5, 2.0);
1115 if ratio > 1.15 {
1116 if allow_growth {
1117 ((units as f64 * ratio).ceil() as u32).clamp(units, units.saturating_mul(2))
1118 } else {
1119 units
1120 }
1121 } else if ratio < 0.85 {
1122 ((units as f64 * ratio).floor() as u32).clamp((units / 2).max(1), units)
1123 } else {
1124 units
1125 }
1126}
1127
1128struct Cursor {
1129 work: Work,
1130 x: u32,
1131 y: u32,
1132 band_height: u32,
1133 step_rows: u32,
1134}
1135
1136impl Cursor {
1137 #[cfg(test)]
1138 fn new(work: Work) -> Result<Self, String> {
1139 Self::stepped(work, 1, 1)
1140 }
1141
1142 fn stepped(work: Work, rows: u32, columns: u32) -> Result<Self, String> {
1143 let rows = rows.max(1);
1144 let columns = columns.max(1);
1145 match work {
1146 Work::Bands { width, height } | Work::Tiles { width, height, .. }
1147 if width == 0 || height == 0 =>
1148 {
1149 return Err("pixel work needs nonzero width and height".into());
1150 }
1151 Work::Tiles { tile_width: 0, .. } => return Err("tile width must be nonzero".into()),
1152 Work::Dispatches { count: 0 } => return Err("dispatch count must be nonzero".into()),
1153 _ => {}
1154 }
1155 let work = match work {
1156 Work::Tiles {
1157 width,
1158 height,
1159 tile_width,
1160 } => Work::Tiles {
1161 width,
1162 height,
1163 tile_width: (tile_width / columns * columns).max(columns),
1164 },
1165 other => other,
1166 };
1167 Ok(Self {
1168 work,
1169 x: 0,
1170 y: 0,
1171 band_height: 0,
1172 step_rows: match work {
1173 Work::Dispatches { .. } => 1,
1174 _ => rows,
1175 },
1176 })
1177 }
1178
1179 fn grant(&self, units: u32) -> u32 {
1180 match self.work {
1181 Work::Bands { .. } | Work::Tiles { .. } => {
1182 (units / self.step_rows * self.step_rows).max(self.step_rows)
1183 }
1184 Work::Dispatches { .. } => units.max(1),
1185 }
1186 }
1187
1188 fn available(&self) -> u32 {
1189 match self.work {
1190 Work::Bands { height, .. } | Work::Tiles { height, .. } => height,
1191 Work::Dispatches { count } => count,
1192 }
1193 }
1194
1195 fn weight(&self, slice: Slice) -> f64 {
1196 match (self.work, slice) {
1197 (
1198 Work::Bands { width, .. },
1199 Slice::Pixels {
1200 width: columns,
1201 height,
1202 ..
1203 },
1204 ) => height as f64 * columns as f64 / width as f64,
1205 (
1206 Work::Tiles {
1207 width, tile_width, ..
1208 },
1209 Slice::Pixels {
1210 width: columns,
1211 height,
1212 ..
1213 },
1214 ) => height as f64 * columns as f64 / tile_width.min(width) as f64,
1215 (_, Slice::Pixels { height, .. }) => height as f64,
1216 (_, Slice::Dispatches { count, .. }) => count as f64,
1217 }
1218 }
1219
1220 fn next(&mut self, units: u32) -> Option<Slice> {
1221 match self.work {
1222 Work::Bands { width, height } => {
1223 if self.y >= height {
1224 return None;
1225 }
1226 let rows = self.grant(units).min(height - self.y);
1227 let slice = Slice::Pixels {
1228 x: 0,
1229 y: self.y,
1230 width,
1231 height: rows,
1232 };
1233 self.y += rows;
1234 Some(slice)
1235 }
1236 Work::Tiles {
1237 width,
1238 height,
1239 tile_width,
1240 } => {
1241 if self.y >= height {
1242 return None;
1243 }
1244 if self.x == 0 {
1245 self.band_height = self.grant(units).min(height - self.y);
1246 }
1247 let columns = tile_width.min(width - self.x);
1248 let slice = Slice::Pixels {
1249 x: self.x,
1250 y: self.y,
1251 width: columns,
1252 height: self.band_height,
1253 };
1254 self.x += columns;
1255 if self.x == width {
1256 self.x = 0;
1257 self.y += self.band_height;
1258 }
1259 Some(slice)
1260 }
1261 Work::Dispatches { count } => {
1262 if self.y >= count {
1263 return None;
1264 }
1265 let size = units.max(1).min(count - self.y);
1266 let slice = Slice::Dispatches {
1267 start: self.y,
1268 count: size,
1269 };
1270 self.y += size;
1271 Some(slice)
1272 }
1273 }
1274 }
1275}
1276
1277struct Timestamp {
1278 slots: [TimestampSlot; 2],
1279 period: f64,
1280}
1281
1282struct TimestampSlot {
1283 queries: wgpu::QuerySet,
1284 resolve: wgpu::Buffer,
1285 read: wgpu::Buffer,
1286}
1287
1288impl Timestamp {
1289 fn new(device: &wgpu::Device, queue: &wgpu::Queue) -> Self {
1290 Self {
1291 slots: std::array::from_fn(|_| TimestampSlot::new(device)),
1292 period: f64::from(queue.get_timestamp_period()),
1293 }
1294 }
1295
1296 fn read_ms(
1297 &self,
1298 device: &wgpu::Device,
1299 slot: usize,
1300 index: wgpu::SubmissionIndex,
1301 ) -> Option<f64> {
1302 let timing = &self.slots[slot];
1303 let (send, receive) = mpsc::channel();
1304 timing
1305 .read
1306 .slice(..)
1307 .map_async(wgpu::MapMode::Read, move |result| {
1308 let _ = send.send(result);
1309 });
1310 device
1311 .poll(wgpu::PollType::Wait {
1312 submission_index: Some(index),
1313 timeout: None,
1314 })
1315 .ok()?;
1316 receive.recv().ok()?.ok()?;
1317 let data = timing.read.slice(..).get_mapped_range().ok()?;
1318 let first = u64::from_le_bytes(data[0..8].try_into().ok()?);
1319 let last = u64::from_le_bytes(data[8..16].try_into().ok()?);
1320 drop(data);
1321 timing.read.unmap();
1322 Some(last.saturating_sub(first) as f64 * self.period / 1_000_000.0)
1323 }
1324}
1325
1326pub struct SliceTimer {
1327 stamp: Timestamp,
1328}
1329
1330impl SliceTimer {
1331 pub fn new(device: &wgpu::Device, queue: &wgpu::Queue) -> Option<Self> {
1332 device
1333 .features()
1334 .contains(
1335 wgpu::Features::TIMESTAMP_QUERY | wgpu::Features::TIMESTAMP_QUERY_INSIDE_ENCODERS,
1336 )
1337 .then(|| Self {
1338 stamp: Timestamp::new(device, queue),
1339 })
1340 }
1341
1342 pub fn begin(&self, encoder: &mut wgpu::CommandEncoder) {
1343 encoder.write_timestamp(&self.stamp.slots[0].queries, 0);
1344 }
1345
1346 pub fn end(&self, encoder: &mut wgpu::CommandEncoder) {
1347 let timing = &self.stamp.slots[0];
1348 encoder.write_timestamp(&timing.queries, 1);
1349 encoder.resolve_query_set(&timing.queries, 0..2, &timing.resolve, 0);
1350 encoder.copy_buffer_to_buffer(&timing.resolve, 0, &timing.read, 0, 16);
1351 }
1352
1353 pub fn read_ms(&self, device: &wgpu::Device, index: wgpu::SubmissionIndex) -> Option<f64> {
1354 self.stamp.read_ms(device, 0, index)
1355 }
1356}
1357
1358impl TimestampSlot {
1359 fn new(device: &wgpu::Device) -> Self {
1360 Self {
1361 queries: device.create_query_set(&wgpu::QuerySetDescriptor {
1362 label: Some("pacer timestamps"),
1363 ty: wgpu::QueryType::Timestamp,
1364 count: 2,
1365 }),
1366 resolve: device.create_buffer(&wgpu::BufferDescriptor {
1367 label: Some("pacer timestamp resolve"),
1368 size: 16,
1369 usage: wgpu::BufferUsages::QUERY_RESOLVE | wgpu::BufferUsages::COPY_SRC,
1370 mapped_at_creation: false,
1371 }),
1372 read: device.create_buffer(&wgpu::BufferDescriptor {
1373 label: Some("pacer timestamp read"),
1374 size: 16,
1375 usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
1376 mapped_at_creation: false,
1377 }),
1378 }
1379 }
1380}
1381
1382#[cfg(test)]
1383mod tests {
1384 use super::*;
1385
1386 #[test]
1387 fn bands_and_tiles_cover_every_pixel_in_order() {
1388 for work in [
1389 Work::Bands {
1390 width: 7,
1391 height: 11,
1392 },
1393 Work::Tiles {
1394 width: 7,
1395 height: 11,
1396 tile_width: 3,
1397 },
1398 ] {
1399 let mut cursor = Cursor::new(work).unwrap();
1400 let mut slices = Vec::new();
1401 let mut pixels = Vec::new();
1402 while let Some(
1403 slice @ Slice::Pixels {
1404 x,
1405 y,
1406 width,
1407 height,
1408 },
1409 ) = cursor.next(4)
1410 {
1411 slices.push(slice);
1412 for row in y..y + height {
1413 for column in x..x + width {
1414 pixels.push((column, row));
1415 }
1416 }
1417 }
1418 assert_eq!(pixels.len(), 77);
1419 pixels.sort_unstable();
1420 let expected: Vec<_> = (0..11).flat_map(|y| (0..7).map(move |x| (x, y))).collect();
1421 let mut expected_sorted = expected.clone();
1422 expected_sorted.sort_unstable();
1423 assert_eq!(pixels, expected_sorted);
1424 let expected_slices = match work {
1425 Work::Bands { .. } => vec![
1426 Slice::Pixels {
1427 x: 0,
1428 y: 0,
1429 width: 7,
1430 height: 4,
1431 },
1432 Slice::Pixels {
1433 x: 0,
1434 y: 4,
1435 width: 7,
1436 height: 4,
1437 },
1438 Slice::Pixels {
1439 x: 0,
1440 y: 8,
1441 width: 7,
1442 height: 3,
1443 },
1444 ],
1445 Work::Tiles { .. } => [
1446 (0, 0, 3, 4),
1447 (3, 0, 3, 4),
1448 (6, 0, 1, 4),
1449 (0, 4, 3, 4),
1450 (3, 4, 3, 4),
1451 (6, 4, 1, 4),
1452 (0, 8, 3, 3),
1453 (3, 8, 3, 3),
1454 (6, 8, 1, 3),
1455 ]
1456 .into_iter()
1457 .map(|(x, y, width, height)| Slice::Pixels {
1458 x,
1459 y,
1460 width,
1461 height,
1462 })
1463 .collect(),
1464 _ => unreachable!(),
1465 };
1466 assert_eq!(slices, expected_slices);
1467 }
1468 }
1469
1470 #[test]
1471 fn dispatch_slices_have_absolute_starts() {
1472 let mut cursor = Cursor::new(Work::Dispatches { count: 10 }).unwrap();
1473 assert_eq!(
1474 cursor.next(4),
1475 Some(Slice::Dispatches { start: 0, count: 4 })
1476 );
1477 assert_eq!(
1478 cursor.next(4),
1479 Some(Slice::Dispatches { start: 4, count: 4 })
1480 );
1481 assert_eq!(
1482 cursor.next(4),
1483 Some(Slice::Dispatches { start: 8, count: 2 })
1484 );
1485 assert_eq!(cursor.next(4), None);
1486 }
1487
1488 #[test]
1489 fn adaptation_converges_and_stays_bounded() {
1490 assert_eq!(Pacer::default().target_ms, 3.5);
1491 assert_eq!(Pacer::new(1.0).target_ms, 2.0);
1492 assert_eq!(Pacer::new(9.0).target_ms, 6.0);
1493 for initial in [1, 256] {
1494 let mut units = initial;
1495 for _ in 0..30 {
1496 let next = adapt(units, 3.5 / 0.1, true);
1497 assert!(next <= units.saturating_mul(2));
1498 assert!(next >= (units / 2).max(1));
1499 units = next;
1500 }
1501 assert!((30..=40).contains(&units), "{units}");
1502 }
1503 }
1504
1505 struct Synthetic<C> {
1506 cost: C,
1507 timing: Timing,
1508 submissions: u32,
1509 units_done: u32,
1510 trace: Vec<(u32, f64)>,
1511 }
1512
1513 impl<C: FnMut(u32, u32, f64) -> f64> Synthetic<C> {
1514 fn new(timing: Timing, cost: C) -> Self {
1515 Self {
1516 cost,
1517 timing,
1518 submissions: 0,
1519 units_done: 0,
1520 trace: Vec::new(),
1521 }
1522 }
1523 }
1524
1525 impl<C: FnMut(u32, u32, f64) -> f64> Lane for Synthetic<C> {
1526 type Ticket = f64;
1527
1528 fn submit(&mut self, _slot: usize, slice: Slice) -> f64 {
1529 let weight = match slice {
1530 Slice::Pixels { width, height, .. } => (width * height) as f64 / 1000.0,
1531 Slice::Dispatches { count, .. } => count as f64,
1532 };
1533 let ms = (self.cost)(self.submissions, self.units_done, weight);
1534 self.submissions += 1;
1535 self.units_done += weight as u32;
1536 self.trace.push((weight as u32, ms));
1537 ms
1538 }
1539
1540 fn finish(&mut self, _slot: usize, ticket: &f64) -> Result<Measured, String> {
1541 Ok(Measured {
1542 milliseconds: *ticket,
1543 timing: self.timing,
1544 reported: true,
1545 })
1546 }
1547 }
1548
1549 const OVERHEAD_MS: f64 = 0.02;
1550 const STEP_NOISE: f64 = 1.25;
1551 const TOTAL: u32 = 60_000;
1552
1553 fn stepping(outliers: bool) -> impl FnMut(u32, u32, f64) -> f64 {
1554 move |submission, done, weight| {
1555 let per_unit = if (20_000..40_000).contains(&done) {
1556 0.2
1557 } else {
1558 0.05
1559 };
1560 let spike = if outliers && submission % 40 == 39 {
1561 3.0
1562 } else {
1563 1.0
1564 };
1565 (OVERHEAD_MS + weight * per_unit) * spike
1566 }
1567 }
1568
1569 fn unsliced_ms() -> f64 {
1570 OVERHEAD_MS + 40_000.0 * 0.05 + 20_000.0 * 0.2
1571 }
1572
1573 fn legacy(work_units: u32, cost: impl FnMut(u32, u32, f64) -> f64) -> Stats {
1574 fn old_adapt(units: u32, elapsed_ms: f64, target_ms: f64) -> u32 {
1575 let ratio = (target_ms / elapsed_ms).clamp(0.5, 2.0);
1576 if ratio > 1.15 {
1577 ((units as f64 * ratio).ceil() as u32).clamp(units, units.saturating_mul(2))
1578 } else if ratio < 0.85 {
1579 ((units as f64 * ratio).floor() as u32).clamp((units / 2).max(1), units)
1580 } else {
1581 units
1582 }
1583 }
1584 let mut lane = Synthetic::new(Timing::Gpu, cost);
1585 let mut cursor = Cursor::new(Work::Dispatches { count: work_units }).unwrap();
1586 let mut stats = Stats::default();
1587 let mut units = 1;
1588 let mut pending: Option<(f64, u32, u32)> = None;
1589 while let Some(slice) = cursor.next(units) {
1590 let observed = match slice {
1591 Slice::Dispatches { count, .. } => count,
1592 _ => unreachable!(),
1593 };
1594 let current = (lane.submit(0, slice), units, observed);
1595 if let Some((ms, used, seen)) = pending.replace(current) {
1596 stats.milliseconds.push(ms);
1597 if seen == used {
1598 units = old_adapt(used, ms, 3.5);
1599 }
1600 }
1601 }
1602 if let Some((ms, ..)) = pending {
1603 stats.milliseconds.push(ms);
1604 }
1605 stats.wall_ms = stats.milliseconds.iter().sum();
1606 stats
1607 }
1608
1609 fn paced(
1610 work_units: u32,
1611 cost: impl FnMut(u32, u32, f64) -> f64,
1612 ) -> (Stats, Synthetic<impl FnMut(u32, u32, f64) -> f64>) {
1613 let mut lane = Synthetic::new(Timing::Gpu, cost);
1614 let mut pacer = Pacer::default();
1615 let mut stats = pacer
1616 .drive(
1617 "t",
1618 &mut lane,
1619 Work::Dispatches { count: work_units },
1620 |_| {},
1621 )
1622 .unwrap();
1623 stats.wall_ms = stats.milliseconds.iter().sum();
1624 (stats, lane)
1625 }
1626
1627 #[test]
1628 fn truncated_slice_teaches_the_pacer() {
1629 let limits = Limits {
1630 target_ms: 3.5,
1631 ceiling_ms: 6.0,
1632 step: 1,
1633 };
1634 let mut tuning = Tuning::new(12);
1635 tuning.learn(12, 6.0, 0.6, 1000, limits);
1636 assert_eq!(tuning.units, 24);
1637 let mut tuning = Tuning::new(24);
1638 tuning.learn(24, 12.0, 4.8, 1000, limits);
1639 assert_eq!(tuning.units, 12);
1640 let mut tuning = Tuning::new(24);
1641 tuning.learn(24, 3.0, 0.1, 1000, limits);
1642 assert_eq!(tuning.units, 24);
1643 }
1644
1645 #[test]
1646 fn learned_units_never_exceed_the_available_work() {
1647 let limits = Limits {
1648 target_ms: 3.5,
1649 ceiling_ms: 6.0,
1650 step: 1,
1651 };
1652 let mut tuning = Tuning::new(12);
1653 tuning.learn(12, 5.0, 0.5, 5, limits);
1654 assert_eq!(tuning.units, 5);
1655 tuning.learn(5, 5.0, f64::NAN, 3, limits);
1656 assert_eq!(tuning.units, 3);
1657 }
1658
1659 #[test]
1660 fn a_smaller_work_item_reteaches_the_pacer() {
1661 let mut pacer = Pacer::default();
1662 let mut big = Synthetic::new(Timing::Gpu, |_, _, weight| weight * 0.05);
1663 pacer
1664 .drive("t", &mut big, Work::Dispatches { count: 4000 }, |_| {})
1665 .unwrap();
1666 let before = pacer.learned_units("t").unwrap();
1667 assert!((60..=80).contains(&before), "{before}");
1668 let mut small = Synthetic::new(Timing::Gpu, |_, _, weight| weight * 0.2);
1669 pacer
1670 .drive("t", &mut small, Work::Dispatches { count: 40 }, |_| {})
1671 .unwrap();
1672 let after = pacer.learned_units("t").unwrap();
1673 assert!(after <= 40, "{after}");
1674 assert!(after < before);
1675 let mut again = Synthetic::new(Timing::Gpu, |_, _, weight| weight * 0.2);
1676 pacer
1677 .drive("t", &mut again, Work::Dispatches { count: 4000 }, |_| {})
1678 .unwrap();
1679 let settled = pacer.learned_units("t").unwrap();
1680 assert!((14..=20).contains(&settled), "{settled}");
1681 }
1682
1683 #[test]
1684 fn tile_weights_discount_the_narrow_edge_tile() {
1685 let cursor = Cursor::new(Work::Tiles {
1686 width: 10,
1687 height: 8,
1688 tile_width: 4,
1689 })
1690 .unwrap();
1691 let full = Slice::Pixels {
1692 x: 0,
1693 y: 0,
1694 width: 4,
1695 height: 8,
1696 };
1697 let edge = Slice::Pixels {
1698 x: 8,
1699 y: 0,
1700 width: 2,
1701 height: 8,
1702 };
1703 assert_eq!(cursor.weight(full), 8.0);
1704 assert_eq!(cursor.weight(edge), 4.0);
1705 }
1706
1707 #[test]
1708 fn the_ceiling_cuts_after_an_outlier_then_grows_back_slowly() {
1709 let mut pacer = Pacer::default();
1710 let mut lane = Synthetic::new(Timing::Gpu, |submission, _, weight| {
1711 let spike = if submission == 60 { 3.0 } else { 1.0 };
1712 weight * 0.05 * spike
1713 });
1714 let stats = pacer
1715 .drive("t", &mut lane, Work::Dispatches { count: 30_000 }, |_| {})
1716 .unwrap();
1717 let ceiling = pacer.ceiling_ms();
1718 let outlier = stats
1719 .milliseconds
1720 .iter()
1721 .position(|ms| *ms > ceiling)
1722 .expect("the spike exceeds the ceiling");
1723 assert_eq!(outlier, 60);
1724 assert!(
1725 stats.milliseconds[outlier + 2] < ceiling,
1726 "{:?}",
1727 &stats.milliseconds[outlier..outlier + 4]
1728 );
1729 assert!(stats.milliseconds[outlier + 2] < 3.5);
1730 for pair in lane.trace[outlier + 2..outlier + 8].windows(2) {
1731 assert!(pair[1].0 as f64 <= (pair[0].0 as f64 * 1.25).ceil() + 1.0);
1732 }
1733 let later = &stats.milliseconds[outlier + 40..];
1734 assert!(
1735 later.iter().all(|ms| *ms < ceiling),
1736 "{:?}",
1737 later.iter().copied().fold(0.0, f64::max)
1738 );
1739 let tail = lane.trace[lane.trace.len() - 10].0;
1740 assert!((55..=80).contains(&tail), "{tail}");
1741 }
1742
1743 #[test]
1744 fn the_ceiling_is_settable_and_never_below_the_target() {
1745 let mut pacer = Pacer::new(4.0);
1746 assert_eq!(pacer.ceiling_ms(), 6.0);
1747 pacer.set_ceiling_ms(8.0);
1748 assert_eq!(pacer.ceiling_ms(), 8.0);
1749 pacer.set_ceiling_ms(1.0);
1750 assert_eq!(pacer.ceiling_ms(), 4.0);
1751 pacer.set_ceiling_ms(f64::NAN);
1752 assert_eq!(pacer.ceiling_ms(), 4.0);
1753 assert_eq!(Pacer::new(6.0).ceiling_ms(), 6.0);
1754 }
1755
1756 #[test]
1757 fn stats_say_how_every_slice_was_timed() {
1758 for (timing, source) in [
1759 (Timing::Gpu, TimingSource::Gpu),
1760 (Timing::Cpu, TimingSource::Cpu),
1761 ] {
1762 let mut lane = Synthetic::new(timing, |_, _, weight| weight * 0.05);
1763 let stats = Pacer::default()
1764 .drive("t", &mut lane, Work::Dispatches { count: 500 }, |_| {})
1765 .unwrap();
1766 assert_eq!(stats.source(), source);
1767 assert_eq!(stats.timings.len(), stats.count());
1768 assert!(stats.timings.iter().all(|seen| *seen == timing));
1769 }
1770 let mixed = Stats {
1771 milliseconds: vec![1.0, 2.0],
1772 timings: vec![Timing::Gpu, Timing::Cpu],
1773 wall_ms: 3.0,
1774 };
1775 assert_eq!(mixed.source(), TimingSource::Mixed);
1776 assert_eq!(Stats::default().source(), TimingSource::Unmeasured);
1777 }
1778
1779 #[test]
1780 fn percentiles_use_the_nearest_rank() {
1781 let stats = Stats {
1782 milliseconds: (1..=100).map(f64::from).collect(),
1783 ..Stats::default()
1784 };
1785 assert_eq!(stats.percentile_ms(99.0), 99.0);
1786 assert_eq!(stats.percentile_ms(100.0), 100.0);
1787 assert_eq!(stats.percentile_ms(0.0), 1.0);
1788 assert_eq!(Stats::default().percentile_ms(99.0), 0.0);
1789 }
1790
1791 #[test]
1792 fn a_steady_workload_behaves_as_before() {
1793 let steady = |_: u32, _: u32, weight: f64| weight * 0.1;
1794 let before = legacy(20_000, steady);
1795 let (after, _) = paced(20_000, steady);
1796 assert_eq!(after.count(), before.count());
1797 for (new, old) in after.milliseconds.iter().zip(&before.milliseconds) {
1798 assert!((new - old).abs() < 1e-9, "{new} {old}");
1799 }
1800 let late = &after.milliseconds[after.count() / 2..after.count() - 1];
1801 assert!(
1802 late.iter().all(|ms| (2.8..=4.0).contains(ms)),
1803 "{:?}",
1804 late.iter().copied().fold(0.0, f64::max)
1805 );
1806 }
1807
1808 #[test]
1809 fn fixed_units_skip_learning_and_the_ceiling() {
1810 let mut pacer = Pacer::default();
1811 pacer.set_fixed_units(Some(50));
1812 let mut lane = Synthetic::new(Timing::Gpu, |_, _, weight| weight * 1.0);
1813 let stats = pacer
1814 .drive("t", &mut lane, Work::Dispatches { count: 200 }, |_| {})
1815 .unwrap();
1816 assert_eq!(stats.count(), 4);
1817 assert_eq!(pacer.learned_units("t"), None);
1818 }
1819
1820 #[test]
1821 #[ignore = "prints the before and after of the synthetic step and outlier workload"]
1822 fn synthetic_step_and_outliers() {
1823 for (name, outliers) in [("steps", false), ("steps+outliers", true)] {
1824 let before = legacy(TOTAL, stepping(outliers));
1825 let (after, _) = paced(TOTAL, stepping(outliers));
1826 for (which, stats) in [("before", &before), ("after", &after)] {
1827 let calm: Vec<f64> = stats
1828 .milliseconds
1829 .iter()
1830 .enumerate()
1831 .filter(|(index, _)| !outliers || index % 40 != 39)
1832 .map(|(_, ms)| *ms)
1833 .collect();
1834 println!(
1835 "{name} {which}: slices {} median {:.2} p99 {:.2} longest {:.2} over6 {} calm_longest {:.2} calm_over6 {} throughput {:.3}",
1836 stats.count(),
1837 stats.median_ms(),
1838 stats.percentile_ms(99.0),
1839 stats.longest_ms(),
1840 stats.milliseconds.iter().filter(|ms| **ms > 6.0).count(),
1841 calm.iter().copied().fold(0.0, f64::max),
1842 calm.iter().filter(|ms| **ms > 6.0).count(),
1843 unsliced_ms() / stats.wall_ms,
1844 );
1845 }
1846 }
1847 }
1848
1849 #[test]
1850 fn the_step_workload_stays_near_the_target_and_under_the_ceiling() {
1851 let (after, lane) = paced(TOTAL, stepping(true));
1852 let before = legacy(TOTAL, stepping(true));
1853 let calm_over = |stats: &Stats| {
1854 stats
1855 .milliseconds
1856 .iter()
1857 .enumerate()
1858 .filter(|(index, ms)| index % 40 != 39 && **ms > 6.0)
1859 .count()
1860 };
1861 assert!(calm_over(&after) <= calm_over(&before));
1862 assert!(calm_over(&after) <= 2);
1863 assert!(after.count() < 2 * before.count());
1864 let over: Vec<usize> = (0..after.count())
1865 .filter(|index| after.milliseconds[*index] > 6.0)
1866 .collect();
1867 for index in &over {
1868 let sizing = lane.trace.get(index + 2).map(|(_, ms)| *ms);
1869 if let Some(ms) = sizing {
1870 let spiked = (index + 2) % 40 == 39;
1871 assert!(spiked || ms < 6.0, "{index} {ms}");
1872 }
1873 }
1874 assert!(unsliced_ms() / after.wall_ms > 0.7);
1875 }
1876
1877 #[test]
1878 fn pacing_sums_work_and_longest_until_a_turn() {
1879 let costs = [2.0, 5.0, 3.0, 4.0, 1.5, 6.5];
1880 let mut lane = Synthetic::new(Timing::Gpu, |submission, _, _| costs[submission as usize]);
1881 let mut pacer = Pacer::default();
1882 pacer.set_fixed_units(Some(1));
1883 let mut seen = Vec::new();
1884 pacer
1885 .drive("turns", &mut lane, Work::Dispatches { count: 4 }, |ms| {
1886 seen.push(ms);
1887 seen.len() == 3
1888 })
1889 .unwrap();
1890 assert_eq!(seen, [2.0, 5.0, 3.0, 4.0]);
1891 assert_eq!(
1892 pacer.pacing(),
1893 Pacing {
1894 work_ms: 4.0,
1895 longest_ms: 4.0,
1896 slices: 1,
1897 untimed: 0,
1898 source: TimingSource::Gpu,
1899 }
1900 );
1901 let mut lane = Synthetic::new(Timing::Cpu, |submission, _, _| {
1902 costs[4 + submission as usize]
1903 });
1904 pacer
1905 .drive("turns", &mut lane, Work::Dispatches { count: 2 }, |_| {})
1906 .unwrap();
1907 assert_eq!(
1908 pacer.take_pacing(),
1909 Pacing {
1910 work_ms: 12.0,
1911 longest_ms: 6.5,
1912 slices: 3,
1913 untimed: 0,
1914 source: TimingSource::Mixed,
1915 }
1916 );
1917 assert_eq!(pacer.pacing(), Pacing::default());
1918 }
1919
1920 #[test]
1921 fn turns_sum_work_keep_the_longest_and_reset() {
1922 let mut turns = Turns::new(50.0, false);
1923 assert!(!turns.add(20.0));
1924 assert!(!turns.add(f64::NAN));
1925 assert!(!turns.add(29.5));
1926 assert_eq!(turns.pending(), (49.5, 29.5));
1927 assert!(turns.add(4.0));
1928 assert_eq!(turns.pending(), (0.0, 0.0));
1929 assert_eq!(turns.taken(), 1);
1930 assert_eq!(turns.time(|| 7), 7);
1931 assert!(turns.pending().0 >= 0.0);
1932 turns.turn();
1933 assert_eq!(turns.taken(), 2);
1934 assert_eq!(
1935 turn_args(53.5, 4.0004),
1936 ["turn", "--work-ms", "53.500", "--longest-ms", "4.000"]
1937 );
1938 }
1939
1940 #[cfg(unix)]
1941 #[test]
1942 fn turns_from_two_threads_never_overlap() {
1943 if std::env::var_os("PFX_TURN_CHILD").is_some() {
1944 return;
1945 }
1946 let scratch = std::path::Path::new(env!("CARGO_MANIFEST_DIR"))
1947 .join("../../tmp")
1948 .join(format!("pace-turns-{}", std::process::id()));
1949 std::fs::create_dir_all(&scratch).unwrap();
1950 let pgpu = scratch.join("pgpu");
1951 std::fs::write(
1952 &pgpu,
1953 "#!/bin/sh\nmkdir \"$TURN_LOG.held\" 2>/dev/null || echo overlap >> \"$TURN_LOG\"\nsleep 0.1\nprintf '%s\\n' \"$*\" >> \"$TURN_LOG\"\nrmdir \"$TURN_LOG.held\" 2>/dev/null\n",
1954 )
1955 .unwrap();
1956 #[cfg(unix)]
1957 {
1958 use std::os::unix::fs::PermissionsExt;
1959 std::fs::set_permissions(&pgpu, std::fs::Permissions::from_mode(0o700)).unwrap();
1960 }
1961 let log = scratch.join("turns");
1962 let path = std::env::var_os("PATH").unwrap_or_default();
1963 let mut paths = vec![scratch.clone()];
1964 paths.extend(std::env::split_paths(&path));
1965 let status = std::process::Command::new(std::env::current_exe().unwrap())
1966 .args(["--exact", "pace::tests::turns_child", "--test-threads", "1"])
1967 .env("PATH", std::env::join_paths(paths).unwrap())
1968 .env("PFX_TURN_CHILD", "1")
1969 .env("TURN_LOG", &log)
1970 .stdout(std::process::Stdio::null())
1971 .status()
1972 .unwrap();
1973 let lines = std::fs::read_to_string(&log).unwrap_or_default();
1974 std::fs::remove_dir_all(&scratch).unwrap();
1975 assert!(status.success());
1976 let mut lines: Vec<_> = lines.lines().collect();
1977 lines.sort();
1978 assert_eq!(
1979 lines,
1980 [
1981 "turn --work-ms 1.000 --longest-ms 1.000",
1982 "turn --work-ms 2.000 --longest-ms 2.000",
1983 ]
1984 );
1985 }
1986
1987 #[cfg(unix)]
1988 #[test]
1989 fn turns_child() {
1990 if std::env::var_os("PFX_TURN_CHILD").is_none() {
1991 return;
1992 }
1993 let together: Vec<_> = [1.0, 2.0]
1994 .into_iter()
1995 .map(|ms| {
1996 std::thread::spawn(move || {
1997 let mut turns = Turns::new(50.0, true);
1998 turns.add(ms);
1999 turns.turn();
2000 })
2001 })
2002 .collect();
2003 for thread in together {
2004 thread.join().unwrap();
2005 }
2006 }
2007
2008 #[test]
2009 fn the_unit_limit_caps_every_slice() {
2010 let mut lane = Synthetic::new(Timing::Gpu, |_, _, weight| weight * 0.01);
2011 let mut pacer = Pacer::default();
2012 pacer.set_max_units(Some(64));
2013 let stats = pacer
2014 .drive("cap", &mut lane, Work::Dispatches { count: 5000 }, |_| {})
2015 .unwrap();
2016 assert!(lane.trace.iter().all(|&(units, _)| units <= 64));
2017 assert_eq!(
2018 lane.trace.iter().map(|&(units, _)| units).sum::<u32>(),
2019 5000
2020 );
2021 assert!(stats.count() >= 5000 / 64);
2022 pacer.set_max_units(Some(0));
2023 let mut lane = Synthetic::new(Timing::Gpu, |_, _, _| 1.0);
2024 pacer
2025 .drive("cap", &mut lane, Work::Dispatches { count: 3 }, |_| {})
2026 .unwrap();
2027 assert!(lane.trace.iter().all(|&(units, _)| units == 1));
2028 }
2029
2030 #[test]
2031 fn a_blocking_run_covers_the_work_and_passes_errors_on() {
2032 let mut pacer = Pacer::default();
2033 let mut slices = Vec::new();
2034 let stats = pacer
2035 .run_blocking(
2036 "blocking",
2037 Work::Dispatches { count: 9 },
2038 |slice| {
2039 slices.push(slice);
2040 Ok(())
2041 },
2042 |_| {},
2043 )
2044 .unwrap();
2045 assert_eq!(stats.source(), TimingSource::Cpu);
2046 assert_eq!(stats.count(), slices.len());
2047 let mut next = 0;
2048 for slice in slices {
2049 let Slice::Dispatches { start, count } = slice else {
2050 unreachable!()
2051 };
2052 assert_eq!(start, next);
2053 next += count;
2054 }
2055 assert_eq!(next, 9);
2056 assert_eq!(pacer.pacing().slices, stats.count());
2057 let failed = pacer.run_blocking(
2058 "blocking",
2059 Work::Dispatches { count: 9 },
2060 |_| Err("lost".to_string()),
2061 |_| {},
2062 );
2063 assert_eq!(failed.err().as_deref(), Some("lost"));
2064 }
2065
2066 #[test]
2067 #[ignore = "needs a GPU; run under pgpu"]
2068 fn paced_dispatches_cover_everything_and_report_their_timing() {
2069 const GROUPS: u32 = 600;
2070 const ROUNDS: u32 = 2000;
2071 let gpu = pollster::block_on(crate::Gpu::headless()).unwrap();
2072 let device = &gpu.device;
2073 let timed = device.features().contains(
2074 wgpu::Features::TIMESTAMP_QUERY | wgpu::Features::TIMESTAMP_QUERY_INSIDE_ENCODERS,
2075 );
2076 let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
2077 label: Some("pacer test"),
2078 source: wgpu::ShaderSource::Wgsl(
2079 format!(
2080 "@group(0) @binding(0) var<storage, read_write> out: array<u32>;
2081 @group(0) @binding(1) var<uniform> start: vec4<u32>;
2082 @compute @workgroup_size(64)
2083 fn main(@builtin(global_invocation_id) id: vec3<u32>) {{
2084 let index = start.x * 64u + id.x;
2085 var a = index;
2086 for (var i = 0u; i < {ROUNDS}u; i = i + 1u) {{
2087 a = a * 1664525u + 1013904223u;
2088 }}
2089 out[index] = a;
2090 }}"
2091 )
2092 .into(),
2093 ),
2094 });
2095 let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
2096 label: Some("pacer test"),
2097 layout: None,
2098 module: &shader,
2099 entry_point: Some("main"),
2100 compilation_options: Default::default(),
2101 cache: None,
2102 });
2103 let bytes = u64::from(GROUPS) * 64 * 4;
2104 let out = device.create_buffer(&wgpu::BufferDescriptor {
2105 label: Some("out"),
2106 size: bytes,
2107 usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
2108 mapped_at_creation: false,
2109 });
2110 let read = device.create_buffer(&wgpu::BufferDescriptor {
2111 label: Some("read"),
2112 size: bytes,
2113 usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
2114 mapped_at_creation: false,
2115 });
2116 let mut pacer = Pacer::default();
2117 let mut slices = Vec::new();
2118 let stats = pacer
2119 .run(
2120 "gpu test",
2121 device,
2122 &gpu.queue,
2123 Work::Dispatches { count: GROUPS },
2124 |encoder, slice| {
2125 let Slice::Dispatches { start, count } = slice else {
2126 unreachable!()
2127 };
2128 slices.push((start, count));
2129 let uniform = device.create_buffer(&wgpu::BufferDescriptor {
2130 label: Some("start"),
2131 size: 16,
2132 usage: wgpu::BufferUsages::UNIFORM,
2133 mapped_at_creation: true,
2134 });
2135 uniform
2136 .slice(..)
2137 .get_mapped_range_mut()
2138 .expect("a mapped range")
2139 .slice(0..4)
2140 .copy_from_slice(&start.to_le_bytes());
2141 uniform.unmap();
2142 let group = device.create_bind_group(&wgpu::BindGroupDescriptor {
2143 label: None,
2144 layout: &pipeline.get_bind_group_layout(0),
2145 entries: &[
2146 wgpu::BindGroupEntry {
2147 binding: 0,
2148 resource: out.as_entire_binding(),
2149 },
2150 wgpu::BindGroupEntry {
2151 binding: 1,
2152 resource: uniform.as_entire_binding(),
2153 },
2154 ],
2155 });
2156 let mut pass = encoder.begin_compute_pass(&Default::default());
2157 pass.set_pipeline(&pipeline);
2158 pass.set_bind_group(0, &group, &[]);
2159 pass.dispatch_workgroups(count, 1, 1);
2160 },
2161 |_| {},
2162 )
2163 .unwrap();
2164 let mut encoder = device.create_command_encoder(&Default::default());
2165 encoder.copy_buffer_to_buffer(&out, 0, &read, 0, bytes);
2166 gpu.queue.submit(Some(encoder.finish()));
2167 read.slice(..).map_async(wgpu::MapMode::Read, |_| {});
2168 device.poll(wgpu::PollType::wait_indefinitely()).unwrap();
2169 let data = read.slice(..).get_mapped_range().expect("a mapped range");
2170 for index in 0..GROUPS * 64 {
2171 let mut a = index;
2172 for _ in 0..ROUNDS {
2173 a = a.wrapping_mul(1664525).wrapping_add(1013904223);
2174 }
2175 let at = index as usize * 4;
2176 assert_eq!(u32::from_le_bytes(data[at..at + 4].try_into().unwrap()), a);
2177 }
2178 assert_eq!(stats.count(), slices.len());
2179 assert_eq!(stats.timings.len(), stats.count());
2180 assert_eq!(
2181 stats.source(),
2182 if timed {
2183 TimingSource::Gpu
2184 } else {
2185 TimingSource::Cpu
2186 }
2187 );
2188 assert!(
2189 stats
2190 .milliseconds
2191 .iter()
2192 .all(|ms| ms.is_finite() && *ms >= 0.0)
2193 );
2194 println!(
2195 "gpu pacer: {} slices, source {:?}, median {:.3} ms, longest {:.3} ms",
2196 stats.count(),
2197 stats.source(),
2198 stats.median_ms(),
2199 stats.longest_ms()
2200 );
2201 }
2202
2203 fn pixel_slices(
2204 cursor: &mut Cursor,
2205 units: impl Fn(usize) -> u32,
2206 ) -> Vec<(u32, u32, u32, u32)> {
2207 let mut slices = Vec::new();
2208 while let Some(Slice::Pixels {
2209 x,
2210 y,
2211 width,
2212 height,
2213 }) = cursor.next(units(slices.len()))
2214 {
2215 slices.push((x, y, width, height));
2216 }
2217 slices
2218 }
2219
2220 fn assert_exact_tiling(slices: &[(u32, u32, u32, u32)], width: u32, height: u32) {
2221 let mut covered = vec![0u8; width as usize * height as usize];
2222 for &(x, y, columns, rows) in slices {
2223 assert!(x + columns <= width && y + rows <= height);
2224 for row in y..y + rows {
2225 let at = row as usize * width as usize;
2226 for cell in &mut covered[at + x as usize..at + (x + columns) as usize] {
2227 *cell += 1;
2228 }
2229 }
2230 }
2231 assert!(covered.iter().all(|count| *count == 1));
2232 }
2233
2234 #[test]
2235 fn steps_of_eight_tile_a_4k_image_exactly() {
2236 for (work, columns) in [
2237 (
2238 Work::Bands {
2239 width: 3840,
2240 height: 2160,
2241 },
2242 8,
2243 ),
2244 (
2245 Work::Tiles {
2246 width: 3840,
2247 height: 2160,
2248 tile_width: 500,
2249 },
2250 8,
2251 ),
2252 ] {
2253 for sizes in [[1u32, 5, 13, 64, 100], [8, 9, 7, 300, 2]] {
2254 let mut cursor = Cursor::stepped(work, 8, columns).unwrap();
2255 let slices = pixel_slices(&mut cursor, |index| sizes[index % sizes.len()]);
2256 assert_exact_tiling(&slices, 3840, 2160);
2257 for &(x, y, width, height) in &slices {
2258 assert_eq!((x % 8, y % 8), (0, 0));
2259 assert_eq!(height % 8, 0);
2260 assert_eq!(width % 8, 0);
2261 }
2262 }
2263 }
2264 }
2265
2266 #[test]
2267 fn the_edge_slices_are_the_partial_ones() {
2268 let mut cursor = Cursor::stepped(
2269 Work::Bands {
2270 width: 3843,
2271 height: 2163,
2272 },
2273 8,
2274 8,
2275 )
2276 .unwrap();
2277 let slices = pixel_slices(&mut cursor, |_| 100);
2278 assert_exact_tiling(&slices, 3843, 2163);
2279 let (last, rest) = slices.split_last().unwrap();
2280 assert_eq!(last.1 + last.3, 2163);
2281 assert_eq!(last.3 % 8, 2163 % 8);
2282 assert!(
2283 rest.iter()
2284 .all(|slice| slice.3 % 8 == 0 && slice.1 % 8 == 0)
2285 );
2286 assert!(rest.iter().all(|slice| slice.2 == 3843));
2287
2288 let mut cursor = Cursor::stepped(
2289 Work::Tiles {
2290 width: 3843,
2291 height: 2163,
2292 tile_width: 500,
2293 },
2294 8,
2295 8,
2296 )
2297 .unwrap();
2298 let slices = pixel_slices(&mut cursor, |_| 100);
2299 assert_exact_tiling(&slices, 3843, 2163);
2300 for &(x, y, width, height) in &slices {
2301 assert_eq!((x % 8, y % 8), (0, 0));
2302 assert!(width % 8 == 0 || x + width == 3843);
2303 assert!(height % 8 == 0 || y + height == 2163);
2304 assert!(width <= 496);
2305 }
2306 assert!(slices.iter().any(|slice| slice.2 % 8 == 3843 % 8));
2307 assert!(slices.last().is_some_and(|slice| slice.3 % 8 == 2163 % 8));
2308 }
2309
2310 #[test]
2311 fn a_step_larger_than_the_image_sends_the_image() {
2312 let mut cursor = Cursor::stepped(
2313 Work::Tiles {
2314 width: 5,
2315 height: 3,
2316 tile_width: 2,
2317 },
2318 8,
2319 8,
2320 )
2321 .unwrap();
2322 let slices = pixel_slices(&mut cursor, |_| 1);
2323 assert_eq!(slices, [(0, 0, 5, 3)]);
2324 }
2325
2326 fn slice_heights(
2327 work: Work,
2328 step: Option<(u32, u32)>,
2329 cost: impl FnMut(u32, u32, f64) -> f64,
2330 ) -> (Stats, Vec<(u32, f64)>, Pacer) {
2331 let mut lane = Synthetic::new(Timing::Gpu, cost);
2332 let mut pacer = Pacer::default();
2333 if let Some((rows, columns)) = step {
2334 pacer.set_step(rows, columns);
2335 }
2336 let mut stats = pacer.drive("t", &mut lane, work, |_| {}).unwrap();
2337 stats.wall_ms = stats.milliseconds.iter().sum();
2338 (stats, lane.trace, pacer)
2339 }
2340
2341 #[test]
2342 fn the_default_step_cuts_bands_and_tiles_like_dispatches() {
2343 let steady = |_: u32, _: u32, weight: f64| weight * 0.1;
2344 for work in [
2345 Work::Bands {
2346 width: 1000,
2347 height: 20_000,
2348 },
2349 Work::Tiles {
2350 width: 1000,
2351 height: 20_000,
2352 tile_width: 1000,
2353 },
2354 ] {
2355 let (_, reference, _) = slice_heights(Work::Dispatches { count: 20_000 }, None, steady);
2356 let (_, plain, pacer) = slice_heights(work, None, steady);
2357 assert_eq!(pacer.step(), (1, 1));
2358 assert_eq!(plain, reference);
2359 let (_, explicit, _) = slice_heights(work, Some((1, 1)), steady);
2360 assert_eq!(explicit, reference);
2361 let (_, stepped_trace, _) = slice_heights(
2362 Work::Bands {
2363 width: 1000,
2364 height: 20_000,
2365 },
2366 None,
2367 stepping(true),
2368 );
2369 let (_, dispatch_trace, _) =
2370 slice_heights(Work::Dispatches { count: 20_000 }, None, stepping(true));
2371 assert_eq!(stepped_trace, dispatch_trace);
2372 }
2373 }
2374
2375 #[test]
2376 fn a_step_changes_nothing_for_dispatches() {
2377 let steady = |_: u32, _: u32, weight: f64| weight * 0.1;
2378 let (_, plain, _) = slice_heights(Work::Dispatches { count: 5000 }, None, steady);
2379 let (_, stepped, _) = slice_heights(Work::Dispatches { count: 5000 }, Some((8, 8)), steady);
2380 assert_eq!(plain, stepped);
2381 }
2382
2383 #[test]
2384 fn the_step_and_outlier_workloads_stay_under_the_ceiling_with_a_step_of_eight() {
2385 let work = Work::Bands {
2386 width: 1000,
2387 height: TOTAL,
2388 };
2389 let (after, trace, pacer) = slice_heights(work, Some((8, 8)), stepping(true));
2390 let (flat, ..) = slice_heights(work, None, stepping(true));
2391 assert!(
2392 trace
2393 .iter()
2394 .take(trace.len() - 1)
2395 .all(|(rows, _)| rows % 8 == 0)
2396 );
2397 let ceiling = pacer.ceiling_ms();
2398 let calm_over = |stats: &Stats| {
2399 stats
2400 .milliseconds
2401 .iter()
2402 .enumerate()
2403 .filter(|(index, ms)| index % 40 != 39 && **ms > ceiling)
2404 .count()
2405 };
2406 assert!(
2407 calm_over(&after) <= calm_over(&flat) + 1,
2408 "{}",
2409 calm_over(&after)
2410 );
2411 assert!(calm_over(&after) <= 2, "{}", calm_over(&after));
2412 for index in 0..after.count() {
2413 if after.milliseconds[index] > ceiling
2414 && let Some((_, ms)) = trace.get(index + 2)
2415 {
2416 assert!((index + 2) % 40 == 39 || *ms < ceiling, "{index} {ms}");
2417 }
2418 }
2419 assert!(unsliced_ms() / after.wall_ms > 0.7);
2420 let (calm, _, _) = slice_heights(work, Some((8, 8)), stepping(false));
2421 assert!(calm.milliseconds.iter().filter(|ms| **ms > ceiling).count() <= 2);
2422 assert!(unsliced_ms() / calm.wall_ms > 0.9);
2423 }
2424
2425 #[test]
2426 fn one_step_alone_over_the_ceiling_is_sent_and_reported() {
2427 let work = Work::Bands {
2428 width: 1000,
2429 height: 200,
2430 };
2431 let (stats, trace, pacer) = slice_heights(work, Some((8, 8)), |_, _, weight| weight);
2432 assert!(trace.iter().all(|(rows, _)| *rows == 8));
2433 assert_eq!(stats.count(), 25);
2434 let pacing = pacer.pacing();
2435 assert_eq!(pacing.slices, 25);
2436 assert_eq!(pacing.longest_ms, 8.0);
2437 assert!(pacing.longest_ms > pacer.ceiling_ms());
2438 assert_eq!(pacer.learned_units("t"), Some(8));
2439 }
2440
2441 #[test]
2442 fn steps_are_never_below_one() {
2443 let mut pacer = Pacer::default();
2444 pacer.set_step(0, 0);
2445 assert_eq!(pacer.step(), (1, 1));
2446 pacer.set_step(8, 16);
2447 assert_eq!(pacer.step(), (8, 16));
2448 }
2449
2450 #[test]
2451 fn learning_lands_on_the_grid() {
2452 let limits = Limits {
2453 target_ms: 3.5,
2454 ceiling_ms: 6.0,
2455 step: 8,
2456 };
2457 let mut tuning = Tuning::new(8);
2458 tuning.learn(8, 8.0, 0.5, 1000, limits);
2459 assert_eq!(tuning.units, 16);
2460 let mut tuning = Tuning::new(16);
2461 tuning.learn(16, 16.0, 9.0, 1000, limits);
2462 assert_eq!(tuning.units, 8);
2463 let mut tuning = Tuning::new(64);
2464 tuning.learn(64, 64.0, 2.0, 1000, limits);
2465 assert_eq!(tuning.units % 8, 0);
2466 }
2467
2468 fn latency_bound(overhead_ms: f64, per_row_ms: f64) -> impl FnMut(u32, u32, f64) -> f64 {
2469 move |_, _, rows| overhead_ms + per_row_ms * rows
2470 }
2471
2472 fn rows_and_ms(trace: &[(u32, f64)]) -> Vec<(u32, f64)> {
2473 trace.to_vec()
2474 }
2475
2476 #[test]
2477 #[ignore = "prints the slices a latency-bound workload takes to reach the target band"]
2478 fn latency_bound_convergence() {
2479 for (name, step) in [("step 8", Some((8, 8))), ("step 1", None)] {
2480 let work = Work::Bands {
2481 width: 1000,
2482 height: 20_000,
2483 };
2484 let (_, trace, _) = slice_heights(work, step, latency_bound(3.0, 0.012));
2485 let in_band = trace
2486 .iter()
2487 .position(|&(_, ms)| (3.0..=4.2).contains(&ms) && ms >= 3.3);
2488 println!(
2489 "{name}: first slice at 3.3 ms or more {in_band:?}, rows {:?}",
2490 trace.iter().take(14).map(|t| t.0).collect::<Vec<_>>()
2491 );
2492 }
2493 }
2494
2495 #[test]
2496 fn a_latency_bound_band_grows_to_the_target() {
2497 let work = Work::Bands {
2498 width: 1000,
2499 height: 20_000,
2500 };
2501 let (stats, trace, pacer) = slice_heights(work, Some((8, 8)), latency_bound(3.0, 0.012));
2502 let trace = rows_and_ms(&trace);
2503 let reached = trace
2504 .iter()
2505 .position(|&(rows, ms)| rows >= 32 && ms >= 3.3)
2506 .expect("the bands reach the target band");
2507 assert!(reached <= 8, "{reached} {:?}", &trace[..12]);
2508 for &(rows, ms) in &trace[reached..trace.len() - 1] {
2509 assert!((3.3..=4.2).contains(&ms), "{rows} {ms}");
2510 assert_eq!(rows % 8, 0);
2511 }
2512 assert!(stats.longest_ms() < pacer.ceiling_ms());
2513 assert!(matches!(pacer.learned_units("t"), Some(32..=56)));
2514 }
2515
2516 #[test]
2517 fn a_latency_bound_band_regrows_after_an_outlier() {
2518 let work = Work::Bands {
2519 width: 1000,
2520 height: 20_000,
2521 };
2522 let mut spiked = latency_bound(3.0, 0.012);
2523 let (stats, trace, pacer) =
2524 slice_heights(work, Some((8, 8)), move |submission, done, rows| {
2525 let spike = if submission == 30 { 3.0 } else { 1.0 };
2526 spiked(submission, done, rows) * spike
2527 });
2528 assert!(stats.milliseconds[30] > pacer.ceiling_ms());
2529 let after = &trace[31..];
2530 let regrown = after
2531 .iter()
2532 .position(|&(rows, _)| rows >= 32)
2533 .expect("the bands grow back");
2534 assert!(regrown <= 12, "{regrown} {:?}", &after[..14]);
2535 for &(_, ms) in &after[2..after.len() - 1] {
2536 assert!(ms < pacer.ceiling_ms());
2537 }
2538 }
2539
2540 #[test]
2541 fn a_latency_bound_band_sizes_to_the_nearest_step() {
2542 let limits = Limits {
2543 target_ms: 3.5,
2544 ceiling_ms: 6.0,
2545 step: 8,
2546 };
2547 let mut tuning = Tuning::new(16);
2548 tuning.history.push(8.0, 3.0 + 0.012 * 8.0);
2549 tuning.learn(16, 16.0, 3.0 + 0.012 * 16.0, 1000, limits);
2550 assert_eq!(tuning.units, 32);
2551 tuning.learn(32, 32.0, 3.0 + 0.012 * 32.0, 1000, limits);
2552 assert_eq!(tuning.units, 40);
2553 tuning.learn(40, 40.0, 3.0 + 0.012 * 40.0, 1000, limits);
2554 assert_eq!(tuning.units, 48);
2555 tuning.learn(48, 48.0, 3.0 + 0.012 * 48.0, 1000, limits);
2556 assert_eq!(tuning.units, 48);
2557 }
2558
2559 #[test]
2560 fn a_single_size_grows_by_a_step_when_one_more_stays_under_the_ceiling() {
2561 let limits = Limits {
2562 target_ms: 3.5,
2563 ceiling_ms: 6.0,
2564 step: 8,
2565 };
2566 let mut tuning = Tuning::new(8);
2567 tuning.learn(8, 8.0, 3.096, 1000, limits);
2568 assert_eq!(tuning.units, 16);
2569 let mut tuning = Tuning::new(16);
2570 tuning.learn(16, 16.0, 3.2, 1000, limits);
2571 assert_eq!(tuning.units, 24);
2572 let mut tuning = Tuning::new(8);
2573 tuning.learn(8, 8.0, 5.0, 1000, limits);
2574 assert_eq!(tuning.units, 8);
2575 let mut tuning = Tuning::new(64);
2576 tuning.learn(64, 64.0, 3.2, 1000, limits);
2577 assert_eq!(tuning.units, 64);
2578 let mut tuning = Tuning::new(8);
2579 tuning.learn(8, 8.0, 3.096, 8, limits);
2580 assert_eq!(tuning.units, 8);
2581 }
2582
2583 #[test]
2584 fn recovery_grows_by_a_step_when_the_cost_is_overhead() {
2585 let limits = Limits {
2586 target_ms: 3.5,
2587 ceiling_ms: 6.0,
2588 step: 8,
2589 };
2590 let mut tuning = Tuning::new(16);
2591 tuning.recovering = RECOVERY_SLICES;
2592 tuning.history.push(8.0, 3.096);
2593 tuning.learn(16, 16.0, 3.192, 1000, limits);
2594 assert!(tuning.units >= 24, "{}", tuning.units);
2595 }
2596
2597 #[test]
2598 fn the_fit_finds_the_overhead_and_shrugs_off_an_outlier() {
2599 let mut history = History::new();
2600 for rows in [8.0, 16.0, 24.0, 16.0, 40.0, 24.0, 32.0] {
2601 history.push(rows, 3.0 + 0.012 * rows);
2602 }
2603 history.push(24.0, 3.0 + 0.012 * 24.0 + 9.0);
2604 let model = history.fit().unwrap();
2605 assert!((model.overhead_ms - 3.0).abs() < 0.05, "{model:?}");
2606 assert!((model.per_unit_ms - 0.012).abs() < 0.002, "{model:?}");
2607 let mut same = History::new();
2608 same.push(8.0, 3.0);
2609 same.push(8.0, 3.1);
2610 same.push(9.0, 3.1);
2611 assert!(same.fit().is_none());
2612 let mut linear = History::new();
2613 for rows in [8.0, 16.0, 32.0] {
2614 linear.push(rows, rows * 0.1);
2615 }
2616 assert!(linear.fit().unwrap().overhead_ms < 1e-9);
2617 }
2618
2619 #[test]
2620 fn a_shift_in_cost_forgets_the_old_fit() {
2621 let mut history = History::new();
2622 for rows in [8.0, 16.0, 24.0, 32.0] {
2623 history.record(rows, 3.0 + 0.012 * rows, false);
2624 }
2625 assert_eq!(history.points().len(), 4);
2626 history.record(32.0, 0.6, false);
2627 assert_eq!(history.points(), [(32.0, 0.6)]);
2628 history.record(32.0, 20.0, true);
2629 assert_eq!(history.points(), [(32.0, 20.0)]);
2630 }
2631
2632 #[test]
2633 fn a_linear_workload_on_a_step_of_eight_stays_within_a_step_of_the_target() {
2634 let work = Work::Bands {
2635 width: 1000,
2636 height: 20_000,
2637 };
2638 let (stats, trace, pacer) = slice_heights(work, Some((8, 8)), |_, _, rows| rows * 0.1);
2639 let tail = &trace[trace.len() / 2..trace.len() - 1];
2640 assert!(
2641 tail.iter()
2642 .all(|&(rows, ms)| matches!(rows, 32 | 40) && ms <= 4.0 + 1e-9)
2643 );
2644 assert!(stats.longest_ms() <= pacer.ceiling_ms());
2645 }
2646
2647 #[test]
2648 fn the_wall_fallback_turns_a_light_job_four_tenths_of_a_second_apart() {
2649 let began = Instant::now();
2650 let mut turns = Turns::new(50.0, false);
2651 turns.last_turn = began;
2652 turns.last_event = began;
2653 let mut last = began;
2654 let mut gaps = Vec::new();
2655 for tick in 1..=500u32 {
2656 let now = began + std::time::Duration::from_millis(u64::from(tick) * 10);
2657 if turns.add_at(0.5, now) {
2658 gaps.push(now.duration_since(last));
2659 last = now;
2660 }
2661 }
2662 assert_eq!(turns.taken(), 12);
2663 assert!(
2664 gaps.iter()
2665 .all(|gap| *gap == std::time::Duration::from_millis(390))
2666 );
2667 let slow = Turns::new(50.0, false).with_wall_ms(1000.0);
2668 assert!(!slow.due_at(slow.last_turn + std::time::Duration::from_millis(999)));
2669 assert!(slow.due_at(slow.last_turn + std::time::Duration::from_millis(1000)));
2670 let kept = Turns::new(50.0, false).with_wall_ms(f64::NAN);
2671 assert_eq!(kept.wall_ms, TURN_WALL_MS);
2672 }
2673
2674 #[test]
2675 #[ignore = "needs a GPU; run under pgpu"]
2676 fn sliced_path_tracer_stand_in_on_the_workgroup_grid() {
2677 const WIDTH: u32 = 3840;
2678 const HEIGHT: u32 = 2160;
2679 let gpu = pollster::block_on(crate::Gpu::headless()).unwrap();
2680 let device = &gpu.device;
2681 for scale in [1u32, 2] {
2682 let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
2683 label: Some("pacer step test"),
2684 source: wgpu::ShaderSource::Wgsl(
2685 "@group(0) @binding(0) var<storage, read_write> out: array<u32>;
2686 @group(0) @binding(1) var<uniform> slice: vec4<u32>;
2687 @group(0) @binding(2) var<uniform> image: vec4<u32>;
2688 @compute @workgroup_size(8, 8)
2689 fn main(@builtin(global_invocation_id) id: vec3<u32>) {
2690 if (id.x >= slice.z || id.y >= slice.w) {
2691 return;
2692 }
2693 let px = slice.x + id.x;
2694 let py = slice.y + id.y;
2695 let rounds = SCALEu * (1200u + (((px / 24u) * 31u + (py / 24u) * 17u) % 61u) * 360u
2696 + select(0u, 18000u, ((px / 96u) + (py / 96u)) % 5u == 0u));
2697 var a = py * image.x + px;
2698 for (var i = 0u; i < rounds; i = i + 1u) {
2699 a = a * 1664525u + 1013904223u;
2700 }
2701 out[py * image.x + px] = a;
2702 }"
2703 .replace("SCALE", &scale.to_string())
2704 .into(),
2705 ),
2706 });
2707 let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
2708 label: Some("pacer step test"),
2709 layout: None,
2710 module: &shader,
2711 entry_point: Some("main"),
2712 compilation_options: Default::default(),
2713 cache: None,
2714 });
2715 let bytes = u64::from(WIDTH) * u64::from(HEIGHT) * 4;
2716 let out = device.create_buffer(&wgpu::BufferDescriptor {
2717 label: Some("out"),
2718 size: bytes,
2719 usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
2720 mapped_at_creation: false,
2721 });
2722 let read = device.create_buffer(&wgpu::BufferDescriptor {
2723 label: Some("read"),
2724 size: bytes,
2725 usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
2726 mapped_at_creation: false,
2727 });
2728 let uniform = |values: [u32; 4]| {
2729 let buffer = device.create_buffer(&wgpu::BufferDescriptor {
2730 label: Some("uniform"),
2731 size: 16,
2732 usage: wgpu::BufferUsages::UNIFORM,
2733 mapped_at_creation: true,
2734 });
2735 for (at, value) in values.iter().enumerate() {
2736 buffer
2737 .slice(..)
2738 .get_mapped_range_mut()
2739 .expect("a mapped range")
2740 .slice(at * 4..at * 4 + 4)
2741 .copy_from_slice(&value.to_le_bytes());
2742 }
2743 buffer.unmap();
2744 buffer
2745 };
2746 let mut turns = Turns::default();
2747 let mut run = |label: &str, step: u32, fixed: Option<u32>| {
2748 let mut pacer = Pacer::default();
2749 pacer.set_step(step, step);
2750 pacer.set_fixed_units(fixed);
2751 let mut heights = Vec::new();
2752 let stats = pacer
2753 .run(
2754 label,
2755 device,
2756 &gpu.queue,
2757 Work::Bands {
2758 width: WIDTH,
2759 height: HEIGHT,
2760 },
2761 |encoder, slice| {
2762 let Slice::Pixels {
2763 x,
2764 y,
2765 width,
2766 height,
2767 } = slice
2768 else {
2769 unreachable!()
2770 };
2771 heights.push((y, height));
2772 let slice_uniform = uniform([x, y, width, height]);
2773 let image_uniform = uniform([WIDTH, HEIGHT, 0, 0]);
2774 let group = device.create_bind_group(&wgpu::BindGroupDescriptor {
2775 label: None,
2776 layout: &pipeline.get_bind_group_layout(0),
2777 entries: &[
2778 wgpu::BindGroupEntry {
2779 binding: 0,
2780 resource: out.as_entire_binding(),
2781 },
2782 wgpu::BindGroupEntry {
2783 binding: 1,
2784 resource: slice_uniform.as_entire_binding(),
2785 },
2786 wgpu::BindGroupEntry {
2787 binding: 2,
2788 resource: image_uniform.as_entire_binding(),
2789 },
2790 ],
2791 });
2792 let mut pass = encoder.begin_compute_pass(&Default::default());
2793 pass.set_pipeline(&pipeline);
2794 pass.set_bind_group(0, &group, &[]);
2795 pass.dispatch_workgroups(width.div_ceil(8), height.div_ceil(8), 1);
2796 },
2797 |ms| turns.add(ms),
2798 )
2799 .unwrap();
2800 let mut encoder = device.create_command_encoder(&Default::default());
2801 encoder.copy_buffer_to_buffer(&out, 0, &read, 0, bytes);
2802 gpu.queue.submit(Some(encoder.finish()));
2803 read.slice(..).map_async(wgpu::MapMode::Read, |_| {});
2804 device.poll(wgpu::PollType::wait_indefinitely()).unwrap();
2805 let data = read
2806 .slice(..)
2807 .get_mapped_range()
2808 .expect("a mapped range")
2809 .to_vec();
2810 read.unmap();
2811 (stats, heights, data)
2812 };
2813 let mut best_unsliced = f64::MAX;
2814 let mut reference = Vec::new();
2815 for round in 0..3 {
2816 let (stats, _, data) = run("step test unsliced", 1, Some(HEIGHT));
2817 assert_eq!(stats.count(), 1);
2818 if round > 0 {
2819 best_unsliced = best_unsliced.min(stats.milliseconds[0]);
2820 }
2821 reference = data;
2822 }
2823 println!("scale {scale}: unsliced {best_unsliced:.3} ms");
2824 let mut ratios = Vec::new();
2825 for step in [1u32, 8] {
2826 let label = format!("step test {step}");
2827 let mut best = f64::MAX;
2828 for round in 0..2 {
2829 let (stats, heights, data) = run(&label, step, None);
2830 assert!(data == reference);
2831 assert!(
2832 stats.source() == TimingSource::Gpu || stats.source() == TimingSource::Cpu
2833 );
2834 for &(y, height) in &heights[..heights.len() - 1] {
2835 assert_eq!((y % step, height % step), (0, 0));
2836 }
2837 let sum: f64 = stats.milliseconds.iter().sum();
2838 println!(
2839 "scale {scale} step {step} run {round}: {} slices, summed {:.3} ms, ratio {:.3}, median {:.3} ms, longest {:.3} ms, source {:?}",
2840 stats.count(),
2841 sum,
2842 sum / best_unsliced,
2843 stats.median_ms(),
2844 stats.longest_ms(),
2845 stats.source()
2846 );
2847 if round > 0 {
2848 best = best.min(sum);
2849 }
2850 }
2851 ratios.push(best / best_unsliced);
2852 }
2853 println!(
2854 "scale {scale} summed GPU time over unsliced: step 1 {:.3}, step 8 {:.3}",
2855 ratios[0], ratios[1]
2856 );
2857 assert!(ratios[1] <= ratios[0] * STEP_NOISE);
2858 }
2859 }
2860
2861 fn moment(base: Instant, ms: u64, frozen_ms: f64) -> Moment {
2862 Moment {
2863 at: base + std::time::Duration::from_millis(ms),
2864 frozen_ms,
2865 }
2866 }
2867
2868 #[test]
2869 fn frozen_time_comes_out_of_every_wall_interval() {
2870 assert_eq!(unfrozen_ms(100.0, 10.0, 70.0), 40.0);
2871 assert_eq!(unfrozen_ms(100.0, 10.0, 500.0), 0.0);
2872 assert_eq!(unfrozen_ms(100.0, 70.0, 10.0), 100.0);
2873 let base = Instant::now();
2874 let after = moment(base, 127_004, 127_000.0);
2875 assert_eq!(after.ms_since(moment(base, 0, 0.0)), 4.0);
2876 assert_eq!(moment(base, 0, 0.0).ms_since(after), 0.0);
2877 let scratch = std::path::Path::new(env!("CARGO_MANIFEST_DIR"))
2878 .join("../../tmp")
2879 .join(format!("pace-frozen-{}", std::process::id()));
2880 std::fs::create_dir_all(&scratch).unwrap();
2881 let file = scratch.join("frozen");
2882 let frozen = Frozen::at(&file);
2883 assert_eq!(frozen.ms(), 0.0);
2884 std::fs::write(&file, "4413000\n").unwrap();
2885 assert_eq!(frozen.ms(), 4_413_000.0);
2886 std::fs::write(&file, "soon").unwrap();
2887 assert_eq!(frozen.ms(), 0.0);
2888 std::fs::write(&file, "-1").unwrap();
2889 assert_eq!(frozen.ms(), 0.0);
2890 std::fs::remove_dir_all(&scratch).unwrap();
2891 assert_eq!(Frozen::none().ms(), 0.0);
2892 }
2893
2894 #[test]
2895 fn a_fence_timed_slice_leaves_out_freezes_and_turns() {
2896 let base = Instant::now();
2897 let submitted = moment(base, 0, 0.0);
2898 assert_eq!(
2899 fence_ms(submitted, None, moment(base, 120_010, 120_000.0)),
2900 10.0
2901 );
2902 let resumed = moment(base, 3_300, 0.0);
2903 assert_eq!(
2904 fence_ms(submitted, Some(resumed), moment(base, 3_305, 0.0)),
2905 5.0
2906 );
2907 assert_eq!(
2908 fence_ms(
2909 moment(base, 50, 0.0),
2910 Some(moment(base, 20, 0.0)),
2911 moment(base, 54, 0.0)
2912 ),
2913 4.0
2914 );
2915 }
2916
2917 #[test]
2918 fn a_timed_wait_reports_its_unfrozen_time_and_nothing_across_a_turn() {
2919 let base = Instant::now();
2920 let mut turns = Turns::new(50.0, false).with_frozen(Frozen::none());
2921 turns.last_turn = base;
2922 assert!(!turns.add_timed(moment(base, 0, 0.0), moment(base, 1, 0.0), false));
2923 turns.last_turn = base + std::time::Duration::from_secs(90);
2924 assert!(!turns.add_timed(moment(base, 1, 0.0), moment(base, 90_005, 89_996.0), false));
2925 assert_eq!(turns.pending(), (9.0, 8.0));
2926 let mut nested = Turns::new(50.0, false)
2927 .with_frozen(Frozen::none())
2928 .with_wall_ms(1e9);
2929 let mut inner = Turns::new(50.0, false);
2930 let value = nested.time(|| {
2931 inner.add(30.0);
2932 inner.turn();
2933 5
2934 });
2935 assert_eq!(value, 5);
2936 assert_eq!(nested.pending(), (0.0, 0.0));
2937 assert_eq!(nested.args(), ["turn"]);
2938 nested.add(2.0);
2939 assert_eq!(
2940 nested.args(),
2941 ["turn", "--work-ms", "2.000", "--longest-ms", "2.000"]
2942 );
2943 nested.turn();
2944 assert_eq!(
2945 nested.args(),
2946 ["turn", "--work-ms", "0.000", "--longest-ms", "0.000"]
2947 );
2948 }
2949
2950 #[test]
2951 fn a_measured_run_reports_gpu_time_and_hands_untimed_slices_on_as_bare() {
2952 let mut pacer = Pacer::default();
2953 pacer.set_fixed_units(Some(1));
2954 let mut seen = Vec::new();
2955 let stats = pacer
2956 .run_measured(
2957 "measured",
2958 Work::Dispatches { count: 4 },
2959 |slice| {
2960 let Slice::Dispatches { start, .. } = slice else {
2961 unreachable!()
2962 };
2963 Ok(match start {
2964 0 => Some(2.5),
2965 1 => None,
2966 2 => Some(f64::NAN),
2967 _ => Some(4.0),
2968 })
2969 },
2970 |ms| seen.push(ms),
2971 )
2972 .unwrap();
2973 assert_eq!(seen.len(), 4);
2974 assert_eq!(seen[0], 2.5);
2975 assert!(seen[1].is_nan() && seen[2].is_nan());
2976 assert_eq!(seen[3], 4.0);
2977 assert_eq!(
2978 stats.timings,
2979 [Timing::Gpu, Timing::Cpu, Timing::Cpu, Timing::Gpu]
2980 );
2981 assert_eq!(stats.milliseconds[0], 2.5);
2982 assert_eq!(stats.milliseconds[3], 4.0);
2983 let pacing = pacer.pacing();
2984 assert_eq!(pacing.work_ms, 6.5);
2985 assert_eq!(pacing.longest_ms, 4.0);
2986 assert_eq!(pacing.slices, 4);
2987 assert_eq!(pacing.untimed, 2);
2988 assert_eq!(pacing.source, TimingSource::Mixed);
2989 let mut turns = Turns::new(50.0, false).with_wall_ms(1e9);
2990 let mut blocking = Pacer::default();
2991 blocking
2992 .run_blocking(
2993 "blocking",
2994 Work::Dispatches { count: 3 },
2995 |_| Ok(()),
2996 |ms| turns.add(ms),
2997 )
2998 .unwrap();
2999 assert_eq!(turns.pending(), (0.0, 0.0));
3000 assert_eq!(turns.args(), ["turn"]);
3001 assert_eq!(blocking.pacing().untimed, blocking.pacing().slices);
3002 assert_eq!(blocking.pacing().work_ms, 0.0);
3003 }
3004
3005 #[cfg(unix)]
3006 #[test]
3007 fn frozen_reports_reach_pgpu_without_the_freeze() {
3008 if std::env::var_os("PFX_TURN_CHILD").is_some() {
3009 return;
3010 }
3011 let scratch = std::path::Path::new(env!("CARGO_MANIFEST_DIR"))
3012 .join("../../tmp")
3013 .join(format!("pace-frozen-turns-{}", std::process::id()));
3014 std::fs::create_dir_all(&scratch).unwrap();
3015 let pgpu = scratch.join("pgpu");
3016 std::fs::write(&pgpu, "#!/bin/sh\nprintf '%s\\n' \"$*\" >> \"$TURN_LOG\"\n").unwrap();
3017 #[cfg(unix)]
3018 {
3019 use std::os::unix::fs::PermissionsExt;
3020 std::fs::set_permissions(&pgpu, std::fs::Permissions::from_mode(0o700)).unwrap();
3021 }
3022 let log = scratch.join("turns");
3023 let path = std::env::var_os("PATH").unwrap_or_default();
3024 let mut paths = vec![scratch.clone()];
3025 paths.extend(std::env::split_paths(&path));
3026 let status = std::process::Command::new(std::env::current_exe().unwrap())
3027 .args([
3028 "--exact",
3029 "pace::tests::frozen_child",
3030 "--test-threads",
3031 "1",
3032 ])
3033 .env("PATH", std::env::join_paths(paths).unwrap())
3034 .env("PFX_TURN_CHILD", "1")
3035 .env("TURN_LOG", &log)
3036 .env_remove(FROZEN_VAR)
3037 .stdout(std::process::Stdio::null())
3038 .status()
3039 .unwrap();
3040 let lines = std::fs::read_to_string(&log).unwrap_or_default();
3041 std::fs::remove_dir_all(&scratch).unwrap();
3042 assert!(status.success());
3043 assert_eq!(
3044 lines.lines().collect::<Vec<_>>(),
3045 [
3046 "turn --work-ms 4.000 --longest-ms 4.000",
3047 "turn --work-ms 0.000 --longest-ms 0.000",
3048 "turn",
3049 "turn",
3050 "turn --work-ms 60.000 --longest-ms 30.000",
3051 ]
3052 );
3053 }
3054
3055 #[cfg(unix)]
3056 #[test]
3057 fn frozen_child() {
3058 if std::env::var_os("PFX_TURN_CHILD").is_none() {
3059 return;
3060 }
3061 let base = Instant::now();
3062 let mut held = Turns::new(50.0, true);
3063 held.last_turn = base;
3064 held.add_timed(moment(base, 0, 0.0), moment(base, 90_004, 90_000.0), false);
3065 assert_eq!(held.taken(), 1);
3066 let mut outer = Turns::new(50.0, true).with_wall_ms(1e9);
3067 let mut inner = Turns::new(50.0, true);
3068 outer.time(|| inner.turn());
3069 outer.turn();
3070 let mut blocking = Turns::new(50.0, true).with_wall_ms(1e9);
3071 Pacer::default()
3072 .run_blocking(
3073 "blocking",
3074 Work::Dispatches { count: 3 },
3075 |_| Ok(()),
3076 |ms| blocking.add(ms),
3077 )
3078 .unwrap();
3079 blocking.turn();
3080 let mut measured = Turns::new(50.0, true).with_wall_ms(1e9);
3081 let mut pacer = Pacer::default();
3082 pacer.set_fixed_units(Some(1));
3083 pacer
3084 .run_measured(
3085 "measured",
3086 Work::Dispatches { count: 2 },
3087 |_| Ok(Some(30.0)),
3088 |ms| measured.add(ms),
3089 )
3090 .unwrap();
3091 assert_eq!(measured.taken(), 1);
3092 }
3093
3094 #[test]
3095 #[ignore = "needs a GPU; run under pgpu"]
3096 fn a_blocking_run_times_its_slices_with_timestamps() {
3097 const GROUPS: u32 = 600;
3098 let gpu = pollster::block_on(crate::Gpu::headless()).unwrap();
3099 let device = &gpu.device;
3100 let compile = Instant::now();
3101 let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
3102 label: Some("measured test"),
3103 source: wgpu::ShaderSource::Wgsl(
3104 "@group(0) @binding(0) var<storage, read_write> out: array<u32>;
3105 @compute @workgroup_size(64)
3106 fn main(@builtin(global_invocation_id) id: vec3<u32>) {
3107 var a = id.x;
3108 for (var i = 0u; i < 4000u; i = i + 1u) {
3109 a = a * 1664525u + 1013904223u;
3110 }
3111 out[id.x] = a;
3112 }"
3113 .into(),
3114 ),
3115 });
3116 let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
3117 label: Some("measured test"),
3118 layout: None,
3119 module: &shader,
3120 entry_point: Some("main"),
3121 compilation_options: Default::default(),
3122 cache: None,
3123 });
3124 let compile_ms = compile.elapsed().as_secs_f64() * 1000.0;
3125 let out = device.create_buffer(&wgpu::BufferDescriptor {
3126 label: Some("out"),
3127 size: u64::from(GROUPS) * 64 * 4,
3128 usage: wgpu::BufferUsages::STORAGE,
3129 mapped_at_creation: false,
3130 });
3131 let group = device.create_bind_group(&wgpu::BindGroupDescriptor {
3132 label: None,
3133 layout: &pipeline.get_bind_group_layout(0),
3134 entries: &[wgpu::BindGroupEntry {
3135 binding: 0,
3136 resource: out.as_entire_binding(),
3137 }],
3138 });
3139 let timer = SliceTimer::new(device, &gpu.queue);
3140 let mut walls = Vec::new();
3141 let mut seen = Vec::new();
3142 let mut pacer = Pacer::default();
3143 let stats = pacer
3144 .run_measured(
3145 "measured test",
3146 Work::Dispatches { count: GROUPS },
3147 |slice| {
3148 let Slice::Dispatches { count, .. } = slice else {
3149 unreachable!()
3150 };
3151 let began = Instant::now();
3152 let mut encoder = device.create_command_encoder(&Default::default());
3153 if let Some(timer) = &timer {
3154 timer.begin(&mut encoder);
3155 }
3156 {
3157 let mut pass = encoder.begin_compute_pass(&Default::default());
3158 pass.set_pipeline(&pipeline);
3159 pass.set_bind_group(0, &group, &[]);
3160 pass.dispatch_workgroups(count, 1, 1);
3161 }
3162 if let Some(timer) = &timer {
3163 timer.end(&mut encoder);
3164 }
3165 let index = gpu.queue.submit(Some(encoder.finish()));
3166 let measured = match &timer {
3167 Some(timer) => timer.read_ms(device, index),
3168 None => {
3169 device
3170 .poll(wgpu::PollType::Wait {
3171 submission_index: Some(index),
3172 timeout: None,
3173 })
3174 .map_err(|error| format!("{error:?}"))?;
3175 None
3176 }
3177 };
3178 walls.push(began.elapsed().as_secs_f64() * 1000.0);
3179 Ok(measured)
3180 },
3181 |ms| seen.push(ms),
3182 )
3183 .unwrap();
3184 let timed = timer.is_some();
3185 assert_eq!(seen.len(), stats.count());
3186 if timed {
3187 assert_eq!(stats.source(), TimingSource::Gpu);
3188 assert_eq!(seen, stats.milliseconds);
3189 assert_eq!(pacer.pacing().untimed, 0);
3190 for (gpu_ms, wall_ms) in stats.milliseconds.iter().zip(&walls) {
3191 assert!(*gpu_ms > 0.0 && gpu_ms <= wall_ms, "{gpu_ms} {wall_ms}");
3192 }
3193 } else {
3194 assert_eq!(stats.source(), TimingSource::Cpu);
3195 assert!(seen.iter().all(|ms| ms.is_nan()));
3196 }
3197 println!(
3198 "measured lane: timestamps {timed}, pipeline created in {compile_ms:.3} ms, first slice {:.3} ms GPU in {:.3} ms wall, {} slices, median {:.3} ms, longest {:.3} ms",
3199 stats.milliseconds[0],
3200 walls[0],
3201 stats.count(),
3202 stats.median_ms(),
3203 stats.longest_ms()
3204 );
3205 }
3206
3207 #[test]
3208 fn fixed_units_read_back_as_set() {
3209 let mut pacer = Pacer::default();
3210 assert_eq!(pacer.fixed_units(), None);
3211 pacer.set_fixed_units(Some(0));
3212 assert_eq!(pacer.fixed_units(), Some(1));
3213 pacer.set_fixed_units(Some(12));
3214 assert_eq!(pacer.fixed_units(), Some(12));
3215 }
3216}