@@ -49,6 +49,23 @@ and this project adheres to
4949 forward (shared by the eager and CUDA paths; the eager step no longer
5050 recomputes the dispersion/damping coefficients every time step).
5151
52+ ### Fixed
53+ - ** Disk-staged boundary saving reconstructed a wrong gradient.** Since the
54+ persistent staging session / non-blocking copy stream (PR #81 ), every
55+ ` storage='disk' ` gradient was wrong: max|disk-gpu|/scale 0.2-0.5 for acoustic
56+ 2-D/3-D and 0.8-5.0 for nvar>1 3-D (Elastic3D, DASMu3D); cpu staging with
57+ ` ring_buffers >= 2 ` was off by 1.4-2.4x on the same equations. Three defects in
58+ ` boundary/runtime.cuh ` : (1) ` prefetch_next_backward_chunk_if_needed ` issued
59+ the next chunk early on ` ring_buffers >= 2 ` , but the synchronous disk path is
60+ pinned to slot 0 whatever ` ring_buffers ` says (and defaults to 3 / 2), so the
61+ early H2D overwrote the chunk still being restored -- the predicate is now the
62+ slot assignment; (2) the synchronous-disk enqueue never got the
63+ ` cudaStreamWaitEvent(compute_ready_) ` write-after-read fence the host-staging
64+ branch has; (3) the two nvar>1 restore readers kept the slot-0 override for
65+ cpu staging with ` ring_buffers >= 2 ` . Regression test
66+ ` test/test_boundary_disk_staging_slot.py ` : disk (default knobs and short
67+ chunks) and cpu ring 2 must be bit-exact against gpu-direct.
68+
5269## [ 0.2.0] - 2026-08-24
5370
5471### Added
@@ -135,21 +152,6 @@ and this project adheres to
135152 support Material grid-card layouts.
136153
137154### Fixed
138- - ** Disk-staged boundary saving reconstructed a wrong gradient.** Since the
139- persistent staging session / non-blocking copy stream (PR #81 ), every
140- ` storage='disk' ` gradient was wrong: max|disk-gpu|/scale 0.2-0.5 for acoustic
141- 2-D/3-D and 0.8-5.0 for nvar>1 3-D (Elastic3D, DASMu3D); cpu staging with
142- ` ring_buffers >= 2 ` was off by 1.4-2.4x on the same equations. Three defects in
143- ` boundary/runtime.cuh ` : (1) ` prefetch_next_backward_chunk_if_needed ` issued
144- the next chunk early on ` ring_buffers >= 2 ` , but the synchronous disk path is
145- pinned to slot 0 whatever ` ring_buffers ` says (and defaults to 3 / 2), so the
146- early H2D overwrote the chunk still being restored -- the predicate is now the
147- slot assignment; (2) the synchronous-disk enqueue never got the
148- ` cudaStreamWaitEvent(compute_ready_) ` write-after-read fence the host-staging
149- branch has; (3) the two nvar>1 restore readers kept the slot-0 override for
150- cpu staging with ` ring_buffers >= 2 ` . Regression test
151- ` test/test_boundary_disk_staging_slot.py ` : disk (default knobs and short
152- chunks) and cpu ring 2 must be bit-exact against gpu-direct.
153155- ** CPML aux writes on a domain-decomposition cut tile.** With the strip
154156 allocation, the per-axis PML compute band reaches columns on a cut face that
155157 carry no slab storage; the unclamped index produced a negative offset and the
0 commit comments