diff --git a/Common/Cpp/Containers/GuardPageBuffer.cpp b/Common/Cpp/Containers/GuardPageBuffer.cpp
new file mode 100644
index 0000000000..4665b76270
--- /dev/null
+++ b/Common/Cpp/Containers/GuardPageBuffer.cpp
@@ -0,0 +1,130 @@
+/*  Guard Page Buffer
+ *
+ *  From: https://github.com/PokemonAutomation/
+ *
+ */
+
+#include <new>
+#include "Common/Cpp/Exceptions.h"
+#include "GuardPageBuffer.h"
+
+#ifdef _WIN32
+#include <Windows.h>
+#else
+#include <sys/mman.h>
+#include <unistd.h>
+#endif
+
+namespace PokemonAutomation{
+
+
+
+//  Return the OS page size in bytes: the unit in which the OS maps memory and
+//  sets access rights. It is 4 KB on x86 Windows and Linux, and 16 KB on
+//  Apple Silicon macOS.
+static size_t page_size(){
+#ifdef _WIN32
+    SYSTEM_INFO info;
+    GetSystemInfo(&info);
+    return info.dwPageSize;
+#else
+    return (size_t)sysconf(_SC_PAGESIZE);
+#endif
+}
+
+
+
+GuardPageMemory::~GuardPageMemory(){
+    release();
+}
+
+//  Take over the pages of "x". Afterwards "x" owns no memory, so its
+//  destructor does nothing.
+GuardPageMemory::GuardPageMemory(GuardPageMemory&& x) noexcept
+    : m_base(x.m_base)
+    , m_mapped_bytes(x.m_mapped_bytes)
+    , m_data(x.m_data)
+{
+    x.m_base = nullptr;
+    x.m_mapped_bytes = 0;
+    x.m_data = nullptr;
+}
+void GuardPageMemory::operator=(GuardPageMemory&& x) noexcept{
+    if (this == &x){
+        return;
+    }
+    release();
+    m_base = x.m_base;
+    m_mapped_bytes = x.m_mapped_bytes;
+    m_data = x.m_data;
+    x.m_base = nullptr;
+    x.m_mapped_bytes = 0;
+    x.m_data = nullptr;
+}
+
+//  Unmap the pages if this object still owns them. The destructor calls it,
+//  and so does the move assignment before it takes the pages of another
+//  object.
+void GuardPageMemory::release() noexcept{
+    if (m_base == nullptr){
+        return;
+    }
+#ifdef _WIN32
+    VirtualFree(m_base, 0, MEM_RELEASE);
+#else
+    munmap(m_base, m_mapped_bytes);
+#endif
+    m_base = nullptr;
+}
+
+//  Map the data pages and the guard page, then place the data so that it ends
+//  at the guard page:
+//  1. Round the data size up to whole pages. Add alignment - 1 bytes first,
+//     so that rounding the start down to "alignment" stays inside the mapping.
+//  2. Map all the pages as read-write (mmap() or VirtualAlloc()). New pages
+//     from the OS are all zero, so the buffer starts as zero bytes.
+//  3. Make the last page no-access (mprotect() with PROT_NONE, or
+//     VirtualProtect() with PAGE_NOACCESS). This is the guard page.
+//  4. Put the start of the data at guard - bytes, rounded down to "alignment".
+//  Example with 4096-byte pages, bytes = 72 and alignment = 4: it maps 1 data
+//  page and 1 guard page, and the data starts at byte 4024 of the first page,
+//  so the data ends at byte 4096, where the guard page starts.
+//  If a step fails, it unmaps what it mapped and throws std::bad_alloc.
+GuardPageMemory::GuardPageMemory(size_t bytes, size_t alignment){
+    if (alignment == 0 || (alignment & (alignment - 1)) != 0){
+        throw InternalProgramError(nullptr, PA_CURRENT_FUNCTION, "Alignment must be a power of two.");
+    }
+
+    size_t page = page_size();
+    size_t data_pages = (bytes + alignment - 1 + page - 1) / page;
+    m_mapped_bytes = (data_pages + 1) * page;
+
+#ifdef _WIN32
+    m_base = VirtualAlloc(nullptr, m_mapped_bytes, MEM_RESERVE | MEM_COMMIT, PAGE_READWRITE);
+    if (m_base == nullptr){
+        throw std::bad_alloc();
+    }
+    char* guard = (char*)m_base + data_pages * page;
+    DWORD old_protect;
+    if (!VirtualProtect(guard, page, PAGE_NOACCESS, &old_protect)){
+        VirtualFree(m_base, 0, MEM_RELEASE);
+        throw std::bad_alloc();
+    }
+#else
+    m_base = mmap(nullptr, m_mapped_bytes, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS, -1, 0);
+    if (m_base == MAP_FAILED){
+        throw std::bad_alloc();
+    }
+    char* guard = (char*)m_base + data_pages * page;
+    if (mprotect(guard, page, PROT_NONE) != 0){
+        munmap(m_base, m_mapped_bytes);
+        throw std::bad_alloc();
+    }
+#endif
+
+    m_data = (void*)((size_t)(guard - bytes) & ~(alignment - 1));
+}
+
+
+
+}
diff --git a/Common/Cpp/Containers/GuardPageBuffer.h b/Common/Cpp/Containers/GuardPageBuffer.h
new file mode 100644
index 0000000000..019581baa2
--- /dev/null
+++ b/Common/Cpp/Containers/GuardPageBuffer.h
@@ -0,0 +1,137 @@
+/*  Guard Page Buffer
+ *
+ *  From: https://github.com/PokemonAutomation/
+ *
+ *  A test-only array that crashes on any access past its end. Unit tests of
+ *  SIMD kernels use it to catch reads and writes past the end of a buffer.
+ *
+ */
+
+#ifndef PokemonAutomation_GuardPageBuffer_H
+#define PokemonAutomation_GuardPageBuffer_H
+
+#include <stddef.h>
+#include <type_traits>
+
+namespace PokemonAutomation{
+
+
+//  The memory behind `GuardPageBuffer`, without an element type.
+//
+//  It maps enough whole pages from the OS for "bytes" bytes (plus the room to
+//  round the start down to "alignment"), and then one more page that it marks
+//  as no-access: the guard page. `data()` points to the first byte of the
+//  data, placed so that the data ends right before the guard page.
+//
+//  The memory starts as all zero bytes, and the destructor unmaps it. It can
+//  be moved but not copied. A moved-from object owns no memory.
+//
+//  Throws std::bad_alloc if the OS cannot map or protect the pages, and
+//  InternalProgramError if "alignment" is not a power of two.
+class GuardPageMemory{
+public:
+    ~GuardPageMemory();
+    GuardPageMemory(GuardPageMemory&& x) noexcept;
+    void operator=(GuardPageMemory&& x) noexcept;
+    GuardPageMemory(const GuardPageMemory&) = delete;
+    void operator=(const GuardPageMemory&) = delete;
+
+public:
+    GuardPageMemory(size_t bytes, size_t alignment = 1);
+
+    //  First byte of the data. The data ends right before the guard page
+    //  when "bytes" is a multiple of "alignment".
+    void* data() const{ return m_data; }
+
+private:
+    void release() noexcept;
+
+private:
+    void* m_base;           //  Start of the mapped pages, or nullptr after a move.
+    size_t m_mapped_bytes;  //  Size of the mapping, guard page included.
+    void* m_data;           //  What `data()` returns.
+};
+
+
+
+//  A fixed-size array for unit tests whose last element sits right before a
+//  guard page: a page of memory that the program is not allowed to read or
+//  write. Any access just past the end of the array crashes the test at once.
+//
+//  Why it exists: SIMD kernels work on whole vectors (for example, 16 floats
+//  with AVX-512) and need special code for the last few elements, the "tail".
+//  A tail that loads or stores a full vector touches memory past the end of
+//  the buffer. With a normal heap buffer like `AlignedVector`, that memory
+//  usually belongs to the heap, so nothing happens. The test only crashes on
+//  the rare run where the buffer happens to end next to an unmapped page.
+//  That is how an over-read in the AVX-512 `AudioChannelMixer` tail crashed
+//  CI only now and then. With a GuardPageBuffer, every such run crashes.
+//
+//  How it works: `GuardPageMemory` maps whole pages and makes the last page
+//  no-access. The array is placed so that its end touches that page.
+//  Example with 4 KB pages and `GuardPageBuffer<float> buffer(18)`:
+//  1. The data needs 18 * 4 = 72 bytes, so it maps 1 data page and 1 guard
+//     page (8 KB in total).
+//  2. `buffer.data()` points 72 bytes before the guard page.
+//  3. `buffer[17]` is the last valid element. Reading `buffer[18]`, or a
+//     16-float vector load at `buffer.data() + 16`, crashes.
+//
+//  Where it is used: the kernel unit tests use it for the input and output
+//  buffers of the kernels under test, for example in
+//  `Kernels_AudioChannelMixer_TestInfra.cpp`. `GuardPageImageRGB32` in
+//  `Tests/GuardPageImage.h` uses it for images.
+//
+//  Limits:
+//  - Only trivial types like float and uint32_t. The elements are not
+//    constructed or destroyed, and they start as zero bytes.
+//  - Only accesses after the end are caught, not accesses before the start.
+//  - By default, `data()` is only aligned to alignof(Object). For a kernel
+//    that needs aligned pointers, pass a larger "alignment", for example
+//    PA_ALIGNMENT (64 bytes) for `fft_abs()`. The start is then rounded
+//    down, which can leave a gap of up to alignment - 1 bytes before the
+//    guard page. Accesses in that gap are not caught, so use sizes that are
+//    a multiple of the alignment when you can.
+//  - Each buffer maps at least 2 pages with a system call, so it is slower
+//    than a heap buffer and only meant for tests.
+//
+//  Throws std::bad_alloc if the OS cannot map or protect the pages, and
+//  InternalProgramError if "alignment" is not a power of two.
+template <typename Object>
+class GuardPageBuffer{
+    static_assert(
+        std::is_trivially_default_constructible_v<Object> &&
+        std::is_trivially_destructible_v<Object>,
+        "GuardPageBuffer only holds trivial types. The elements start as zero bytes."
+    );
+
+public:
+    //  Map "size" elements, all zero bytes. "alignment" is the alignment of
+    //  `data()` in bytes and must be a power of two. Example:
+    //  `GuardPageBuffer<float> real(length, PA_ALIGNMENT)` gives `fft_abs()`
+    //  the 64-byte aligned input that it requires.
+    GuardPageBuffer(size_t size, size_t alignment = alignof(Object))
+        : m_memory(size * sizeof(Object), alignment)
+        , m_size(size)
+    {}
+
+    //  Number of elements, as given to the constructor.
+    size_t size() const{ return m_size; }
+
+    //  Element access without a bounds check. An index at or just past
+    //  size() reads the guard page and crashes, which is what tests want.
+    const Object& operator[](size_t index) const{ return data()[index]; }
+          Object& operator[](size_t index)      { return data()[index]; }
+
+    //  Pointer to the first element, to give to the kernel under test.
+    const Object* data() const{ return (const Object*)m_memory.data(); }
+          Object* data()      { return (Object*)m_memory.data(); }
+
+private:
+    GuardPageMemory m_memory;
+    size_t m_size;
+};
+
+
+
+}
+#endif
diff --git a/SerialPrograms/Source/Kernels/AbsFFT/Kernels_AbsFFT_TestInfra.cpp b/SerialPrograms/Source/Kernels/AbsFFT/Kernels_AbsFFT_TestInfra.cpp
index dfa43dffd5..3669498a01 100644
--- a/SerialPrograms/Source/Kernels/AbsFFT/Kernels_AbsFFT_TestInfra.cpp
+++ b/SerialPrograms/Source/Kernels/AbsFFT/Kernels_AbsFFT_TestInfra.cpp
@@ -11,7 +11,7 @@
 #include <cstdio>
 #include <random>
 #include "Common/Compiler.h"
-#include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Kernels_AbsFFT_TestInfra.h"
 
 namespace PokemonAutomation{
@@ -96,11 +96,16 @@ UnitTestResult Test::run(Logger& logger, CancellableScope& scope) const{
     //  Random input with a fixed seed. fft_abs() is destructive on "real",
     //  so each run gets its own copy of the same data.
     //
-    AlignedVector<float> real0(length);
-    AlignedVector<float> real1(length);
-    AlignedVector<float> abs0(half_length);
-    AlignedVector<float> abs1(half_length);
-    AlignedVector<float> weights(length);
+    //  Every buffer ends right before a guard page, so a read or write past
+    //  the end faults on every run. "real" and "abs" must be aligned to 64
+    //  bytes. Their sizes are powers of two, so for all but the smallest
+    //  transforms they still end exactly at the guard page.
+    //
+    GuardPageBuffer<float> real0(length, PA_ALIGNMENT);
+    GuardPageBuffer<float> real1(length, PA_ALIGNMENT);
+    GuardPageBuffer<float> abs0(half_length, PA_ALIGNMENT);
+    GuardPageBuffer<float> abs1(half_length, PA_ALIGNMENT);
+    GuardPageBuffer<float> weights(length);
 
     std::mt19937 generator(42);
     std::uniform_int_distribution<int> sample_distribution(-10000, 10000);
diff --git a/SerialPrograms/Source/Kernels/AudioChannelInterleave/Kernels_AudioChannelInterleave_TestInfra.cpp b/SerialPrograms/Source/Kernels/AudioChannelInterleave/Kernels_AudioChannelInterleave_TestInfra.cpp
index 920f79907f..9064a2ba63 100644
--- a/SerialPrograms/Source/Kernels/AudioChannelInterleave/Kernels_AudioChannelInterleave_TestInfra.cpp
+++ b/SerialPrograms/Source/Kernels/AudioChannelInterleave/Kernels_AudioChannelInterleave_TestInfra.cpp
@@ -5,7 +5,7 @@
  */
 
 #include <random>
-#include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Kernels_AudioChannelInterleave_TestInfra.h"
 
 namespace PokemonAutomation{
@@ -32,9 +32,11 @@ Test::Test(
 UnitTestResult Test::run(Logger& logger, CancellableScope& scope) const{
     size_t total = m_channels * m_samples;
 
-    AlignedVector<float> interleaved0(total);
-    AlignedVector<float> interleaved1(total);
-    AlignedVector<float> strided(total);
+    //  Every buffer ends right before a guard page, so a kernel that reads or
+    //  writes past the end faults on every run.
+    GuardPageBuffer<float> interleaved0(total);
+    GuardPageBuffer<float> interleaved1(total);
+    GuardPageBuffer<float> strided(total);
 
     std::mt19937 generator(42);
 
diff --git a/SerialPrograms/Source/Kernels/AudioChannelMixer/Kernels_AudioChannelMixer_TestInfra.cpp b/SerialPrograms/Source/Kernels/AudioChannelMixer/Kernels_AudioChannelMixer_TestInfra.cpp
index b2734764d3..05e8d31358 100644
--- a/SerialPrograms/Source/Kernels/AudioChannelMixer/Kernels_AudioChannelMixer_TestInfra.cpp
+++ b/SerialPrograms/Source/Kernels/AudioChannelMixer/Kernels_AudioChannelMixer_TestInfra.cpp
@@ -5,7 +5,7 @@
  */
 
 #include <random>
-#include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Kernels_AudioChannelMixer_TestInfra.h"
 
 namespace PokemonAutomation{
@@ -35,9 +35,13 @@ Test::Test(
 
 UnitTestResult Test::run(Logger& logger, CancellableScope& scope) const{
     std::vector<float> matrix(m_in_channels * m_out_channels);
-    AlignedVector<float> in(m_samples * m_in_channels);
-    AlignedVector<float> out0(m_samples * m_out_channels);
-    AlignedVector<float> out1(m_samples * m_out_channels);
+
+    //  The input and both outputs end right before a guard page. A tail that
+    //  reads or writes past the end then faults on every run, not just when
+    //  the heap happens to end there.
+    GuardPageBuffer<float> in(m_samples * m_in_channels);
+    GuardPageBuffer<float> out0(m_samples * m_out_channels);
+    GuardPageBuffer<float> out1(m_samples * m_out_channels);
 
     std::mt19937 generator(42);
 
@@ -63,7 +67,7 @@ UnitTestResult Test::run(Logger& logger, CancellableScope& scope) const{
     );
 
     //  Check for errors.
-    for (size_t c = 0; c < m_samples; c++){
+    for (size_t c = 0; c < m_samples * m_out_channels; c++){
         if (out0[c] != out1[c]){
             return "Mismatch at " + std::to_string(c) + ": " +
                 std::to_string(out0[c]) + " != " + std::to_string(out1[c]);
diff --git a/SerialPrograms/Source/Kernels/AudioChannelMixer/Kernels_AudioChannelMixer_x64_AVX512.cpp b/SerialPrograms/Source/Kernels/AudioChannelMixer/Kernels_AudioChannelMixer_x64_AVX512.cpp
index 5c93de16d0..ef2709fd52 100644
--- a/SerialPrograms/Source/Kernels/AudioChannelMixer/Kernels_AudioChannelMixer_x64_AVX512.cpp
+++ b/SerialPrograms/Source/Kernels/AudioChannelMixer/Kernels_AudioChannelMixer_x64_AVX512.cpp
@@ -126,8 +126,8 @@ void mix_channels_x64_AVX512<2, 1>(
     }
     {
         __mmask16 mask = ((uint32_t)1 << samples) - 1;
-        __m512 out0 = _mm512_mul_ps(i0o0, _mm512_maskz_loadu_ps(mask, in + in_stride*0));
-        out0 = _mm512_fmadd_ps(i1o0, _mm512_loadu_ps(in + in_stride * 1), out0);
+        __m512 out0 = _mm512_mul_ps(i0o0, _mm512_maskz_loadu_ps(mask, in + in_stride * 0));
+        out0 = _mm512_fmadd_ps(i1o0, _mm512_maskz_loadu_ps(mask, in + in_stride * 1), out0);
         _mm512_mask_storeu_ps(out, mask, out0);
     }
 }
diff --git a/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_AudioPipeline_RoundTrip_TestInfra.cpp b/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_AudioPipeline_RoundTrip_TestInfra.cpp
index 5a1be5ee23..e1b77b0098 100644
--- a/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_AudioPipeline_RoundTrip_TestInfra.cpp
+++ b/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_AudioPipeline_RoundTrip_TestInfra.cpp
@@ -10,6 +10,7 @@
 #include "Common/Cpp/Exceptions.h"
 #include "Common/Cpp/PrettyPrint.h"
 #include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Kernels_AudioPipeline_RoundTrip_TestInfra.h"
 #include "Kernels_RawAudio2WorkingAudio.h"
 #include "Kernels_WorkingAudio2RawAudio.h"
@@ -82,8 +83,9 @@ UnitTestResult run_round_trip(
     size_t samples,
     float multiplier
 ){
-    //  Generate random input.
-    AlignedVector<Type> input(samples * channels);
+    //  Generate random input. It ends right before a guard page, so a read
+    //  past the end faults.
+    GuardPageBuffer<Type> input(samples * channels);
     std::mt19937 generator(42);
     if constexpr (std::is_same_v<Type, float>){
         //  Random integer values, all exactly representable.
diff --git a/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_AudioPipeline_Tests.cpp b/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_AudioPipeline_Tests.cpp
index e87632653d..bdc3d1fd45 100644
--- a/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_AudioPipeline_Tests.cpp
+++ b/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_AudioPipeline_Tests.cpp
@@ -9,6 +9,7 @@
 #include "Common/Cpp/Exceptions.h"
 #include "Common/Cpp/PrettyPrint.h"
 #include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Common/Cpp/TestRunners/UnitTestDatabase.h"
 #include "Kernels_AudioPipeline.h"
 #include "Kernels_RawAudio2WorkingAudio.h"
@@ -87,8 +88,9 @@ class RecordingRawFloatListener : public RawAudioStreamListener{
 
 
 
-AlignedVector<float> random_float_input(size_t total_samples, uint32_t seed){
-    AlignedVector<float> input(total_samples);
+//  The buffer ends right before a guard page, so a read past the end faults.
+GuardPageBuffer<float> random_float_input(size_t total_samples, uint32_t seed){
+    GuardPageBuffer<float> input(total_samples);
     std::mt19937 generator(seed);
     //  Integer values, all exactly representable in float.
     std::uniform_int_distribution<long long> distribution(-128, 128);
@@ -125,7 +127,7 @@ class RawFloatGainTest : public UnitTest{
         const size_t sample_rate = 44100;
         const size_t buffer_size = 256; //  Small buffer so the sweep hits multiple chunks.
 
-        AlignedVector<float> input = random_float_input(m_samples * m_channels, 42);
+        GuardPageBuffer<float> input = random_float_input(m_samples * m_channels, 42);
         RawAudio2WorkingAudio raw2working(
             SampleFormat::FLOAT32, sample_rate, m_channels, m_multiplier, buffer_size
         );
@@ -182,7 +184,7 @@ class WorkingFloatGainTest : public UnitTest{
         const size_t buffer_size = 256; //  Small buffer so the sweep hits multiple chunks.
 
         //  Working format: channel-strided, one block per channel.
-        AlignedVector<float> input = random_float_input(m_samples * m_channels, 42);
+        GuardPageBuffer<float> input = random_float_input(m_samples * m_channels, 42);
         WorkingAudio2RawAudio working2raw(
             SampleFormat::FLOAT32, sample_rate, m_channels, m_multiplier, buffer_size
         );
@@ -237,7 +239,10 @@ class FloatGainLiveUpdateTest : public UnitTest{
         const size_t buffer_size = 256;
 
         std::vector<float> gains = {1.0f, 0.5f, 1.0f};
-        std::vector<AlignedVector<float>> inputs(3, random_float_input(batch, 42));
+        std::vector<GuardPageBuffer<float>> inputs;
+        for (size_t b = 0; b < gains.size(); b++){
+            inputs.emplace_back(random_float_input(batch, 42));
+        }
 
         //  RawAudio2WorkingAudio.
         {
diff --git a/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_ChannelMixer_TestInfra.cpp b/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_ChannelMixer_TestInfra.cpp
index 24fd63e9e5..0271ebe848 100644
--- a/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_ChannelMixer_TestInfra.cpp
+++ b/SerialPrograms/Source/Kernels/AudioPipeline/Kernels_ChannelMixer_TestInfra.cpp
@@ -6,6 +6,7 @@
 
 #include <random>
 #include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Kernels_ChannelMixer_TestInfra.h"
 #include "Kernels_ChannelMixer.h"
 
@@ -58,9 +59,10 @@ class RecordingWorkingAudioListener : public WorkingAudioStreamListener{
 
 
 
-//  Random integer samples, all exactly representable in a float.
-AlignedVector<float> make_random_input(size_t channels, size_t samples){
-    AlignedVector<float> input(channels * samples);
+//  Random integer samples, all exactly representable in a float. The buffer
+//  ends right before a guard page, so a read past the end faults.
+GuardPageBuffer<float> make_random_input(size_t channels, size_t samples){
+    GuardPageBuffer<float> input(channels * samples);
     std::mt19937 generator(42);
     std::uniform_int_distribution<long long> distribution(-128, 128);
     for (size_t c = 0; c < input.size(); c++){
@@ -90,7 +92,7 @@ UnitTestResult MixerIdentityTest::run(Logger& logger, CancellableScope& scope) c
         matrix[c + m_channels * c] = 1.0f;
     }
 
-    AlignedVector<float> input = make_random_input(m_channels, m_samples);
+    GuardPageBuffer<float> input = make_random_input(m_channels, m_samples);
 
     const size_t sample_rate = 44100;
 
@@ -135,7 +137,7 @@ UnitTestResult MixerRoundTripTest::run(Logger& logger, CancellableScope& scope)
     //  2 -> 1: Average the two channels.
     std::vector<float> average_matrix{0.5f, 0.5f};
 
-    AlignedVector<float> input = make_random_input(1, m_samples);
+    GuardPageBuffer<float> input = make_random_input(1, m_samples);
 
     const size_t sample_rate = 44100;
 
diff --git a/SerialPrograms/Source/Kernels/AudioResampler/Kernels_AudioResampler_Tests.cpp b/SerialPrograms/Source/Kernels/AudioResampler/Kernels_AudioResampler_Tests.cpp
index e77e5827fa..bc4dc7445c 100644
--- a/SerialPrograms/Source/Kernels/AudioResampler/Kernels_AudioResampler_Tests.cpp
+++ b/SerialPrograms/Source/Kernels/AudioResampler/Kernels_AudioResampler_Tests.cpp
@@ -13,6 +13,7 @@
 #include <random>
 #include "Common/Cpp/CpuId/CpuId.h"
 #include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Common/Cpp/TestRunners/UnitTestDatabase.h"
 #include "Kernels_AudioResampler_BlockKernels.h"
 #include "Kernels_AudioResampler_Tests.h"
@@ -143,8 +144,9 @@ UnitTestResult run_resampler_stream(
     size_t in_rate, size_t out_rate, size_t channels,
     size_t samples, size_t push, uint32_t taps, double beta
 ){
-    //  Generate random input.
-    AlignedVector<float> input(samples * channels);
+    //  Generate random input. It ends right before a guard page, so a read
+    //  past the end faults.
+    GuardPageBuffer<float> input(samples * channels);
     std::mt19937 generator(42);
     std::uniform_real_distribution<float> distribution(-1.f, 1.f);
     for (size_t c = 0; c < input.size(); c++){
@@ -282,14 +284,15 @@ UnitTestResult BlockKernelTest::run(Logger& logger, CancellableScope& scope) con
     //  sample i, so the kernel's in_index is 0 (k = 0 periods in).
     //  The index columns keep counting past the period, so a block that
     //  wraps the period reads samples up to about period_in + block + taps.
-    AlignedVector<float> input(table.period_in() + m_block + m_taps + 128);
+    //  The input and the output both end right before a guard page.
+    GuardPageBuffer<float> input(table.period_in() + m_block + m_taps + 128);
     std::mt19937 generator(42);
     std::uniform_real_distribution<float> distribution(-1.0f, 1.0f);
     for (size_t i = 0; i < input.size(); i++){
         input[i] = distribution(generator);
     }
 
-    AlignedVector<float> got(m_block);
+    GuardPageBuffer<float> got(m_block);
     m_kernel_fn(table, 0, input.data(), m_out_index, got.data());
 
     AlignedVector<float> want = reference_lanes(table, input.data(), m_out_index, m_block);
@@ -341,16 +344,17 @@ UnitTestResult BlockDispatchTest::run(Logger& logger, CancellableScope& scope) c
     );
 
     //  The kernel's input reads may extend sample_block - 1 samples past
-    //  the last window's start.
+    //  the last window's start. The input and the output both end right
+    //  before a guard page.
     const uint32_t* index = table.in_sample_index();
-    AlignedVector<float> input(index[bk.sample_block - 1] + m_taps + bk.sample_block + 64);
+    GuardPageBuffer<float> input(index[bk.sample_block - 1] + m_taps + bk.sample_block + 64);
     std::mt19937 generator(42);
     std::uniform_real_distribution<float> distribution(-1.0f, 1.0f);
     for (size_t i = 0; i < input.size(); i++){
         input[i] = distribution(generator);
     }
 
-    AlignedVector<float> got(bk.sample_block);
+    GuardPageBuffer<float> got(bk.sample_block);
     bk.fp(table, 0, input.data(), 0, got.data());
 
     AlignedVector<float> want = reference_lanes(table, input.data(), 0, bk.sample_block);
diff --git a/SerialPrograms/Source/Kernels/AudioStreamConverter/Kernels_AudioStreamConverter_TestInfra.cpp b/SerialPrograms/Source/Kernels/AudioStreamConverter/Kernels_AudioStreamConverter_TestInfra.cpp
index 83fecdbf8e..7871a8ffd4 100644
--- a/SerialPrograms/Source/Kernels/AudioStreamConverter/Kernels_AudioStreamConverter_TestInfra.cpp
+++ b/SerialPrograms/Source/Kernels/AudioStreamConverter/Kernels_AudioStreamConverter_TestInfra.cpp
@@ -8,7 +8,7 @@
 #include <random>
 #include "Common/Cpp/Exceptions.h"
 #include "Common/Cpp/PrettyPrint.h"
-#include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Kernels_AudioStreamConverter_TestInfra.h"
 
 //  REMOVE
@@ -46,9 +46,11 @@ UnitTestResult Test<Type>::run(Logger& logger, CancellableScope& scope) const{
         throw InternalProgramError(&logger, PA_CURRENT_FUNCTION, "Test size is too small.");
     }
 
-    AlignedVector<Type> ext0(m_samples);
-    AlignedVector<Type> ext1(m_samples);
-    AlignedVector<float> native(m_samples);
+    //  Every buffer ends right before a guard page, so a kernel that reads or
+    //  writes past the end faults on every run.
+    GuardPageBuffer<Type> ext0(m_samples);
+    GuardPageBuffer<Type> ext1(m_samples);
+    GuardPageBuffer<float> native(m_samples);
 
     std::mt19937 generator(42);
 
diff --git a/SerialPrograms/Source/Kernels/BinaryMatrix/Kernels_BinaryMatrix_Tests.cpp b/SerialPrograms/Source/Kernels/BinaryMatrix/Kernels_BinaryMatrix_Tests.cpp
index 22663353f3..9729ba729e 100644
--- a/SerialPrograms/Source/Kernels/BinaryMatrix/Kernels_BinaryMatrix_Tests.cpp
+++ b/SerialPrograms/Source/Kernels/BinaryMatrix/Kernels_BinaryMatrix_Tests.cpp
@@ -17,6 +17,7 @@
 #endif
 //#include "Kernels/BinaryMatrix/Kernels_BinaryMatrix_Arch_64xH_Default.h"
 #include "Kernels/BinaryImageFilters/Kernels_BinaryImage_BasicFilters.h"
+#include "Tests/GuardPageImage.h"
 #include "Kernels_BinaryMatrix_Tests.h"
 
 #include <functional>
@@ -348,7 +349,10 @@ class Test_BinaryMatrix : public UnitTest{
     {}
 
     virtual UnitTestResult run(Logger& logger, CancellableScope& scope) const override{
-        ImageRGB32 image(m_image);
+        //  Copy the test image into guard-page memory with packed rows, so a
+        //  kernel that reads or writes past the last pixel faults.
+        ImageRGB32 file_image(m_image);
+        GuardPageImageRGB32 image(file_image);
 
         if (test_binary_matrix_tile() != 0){
             return false;
diff --git a/SerialPrograms/Source/Kernels/ImageFilters/Kernels_ImageFilter_Tests.cpp b/SerialPrograms/Source/Kernels/ImageFilters/Kernels_ImageFilter_Tests.cpp
index b75ea4273d..6757ab2420 100644
--- a/SerialPrograms/Source/Kernels/ImageFilters/Kernels_ImageFilter_Tests.cpp
+++ b/SerialPrograms/Source/Kernels/ImageFilters/Kernels_ImageFilter_Tests.cpp
@@ -21,6 +21,7 @@
 #include "Kernels/ImageFilters/RGB32_Range/Kernels_ImageFilter_RGB32_Range.h"
 #include "Kernels/ImageFilters/RGB32_EuclideanDistance/Kernels_ImageFilter_RGB32_Euclidean.h"
 #include "Kernels_ImageFilter_Tests.h"
+#include "Tests/GuardPageImage.h"
 #include "Tests/TestUtils.h"
 
 #include <iostream>
@@ -41,7 +42,10 @@ class Test_ImageFilterRGB32Range : public UnitTest{
     {}
 
     virtual UnitTestResult run(Logger& logger, CancellableScope& scope) const override{
-        ImageRGB32 image(m_image);
+        //  Copy the test image into guard-page memory with packed rows, so a
+        //  kernel that reads or writes past the last pixel faults.
+        ImageRGB32 file_image(m_image);
+        GuardPageImageRGB32 image(file_image);
 
         const size_t width = image.width(), height = image.height();
         cout << "Testing filter_rgb32_range(), image size " << width << " x " << height << endl;
@@ -53,8 +57,8 @@ class Test_ImageFilterRGB32Range : public UnitTest{
         const uint32_t mins = uint32_t(min_color);
         const uint32_t maxs = uint32_t(max_color);
 
-        ImageRGB32 image_out(image.width(), image.height());
-        ImageRGB32 image_out_2(image.width(), image.height());
+        GuardPageImageRGB32 image_out(image.width(), image.height());
+        GuardPageImageRGB32 image_out_2(image.width(), image.height());
         size_t pixels_in_range = 0;
 
         const bool replace_color_within_range = true;
@@ -159,7 +163,10 @@ class Test_ImageFilterRGB32Euclidean : public UnitTest{
     {}
 
     virtual UnitTestResult run(Logger& logger, CancellableScope& scope) const override{
-        ImageRGB32 image(m_image);
+        //  Copy the test image into guard-page memory with packed rows, so a
+        //  kernel that reads or writes past the last pixel faults.
+        ImageRGB32 file_image(m_image);
+        GuardPageImageRGB32 image(file_image);
 
         const size_t width = image.width();
         const size_t height = image.height();
@@ -171,8 +178,8 @@ class Test_ImageFilterRGB32Euclidean : public UnitTest{
         double max_dist = 50.0;
         size_t max_dist2 = size_t(max_dist * max_dist);
 
-        ImageRGB32 image_out(image.width(), image.height());
-        ImageRGB32 image_out_2(image.width(), image.height());
+        GuardPageImageRGB32 image_out(image.width(), image.height());
+        GuardPageImageRGB32 image_out_2(image.width(), image.height());
         size_t pixels_in_range = 0;
 
         const bool replace_color_within_range = true;
@@ -286,7 +293,10 @@ class Test_ImageToBlackWhiteRGB32Range : public UnitTest{
     {}
 
     virtual UnitTestResult run(Logger& logger, CancellableScope& scope) const override{
-        ImageRGB32 image(m_image);
+        //  Copy the test image into guard-page memory with packed rows, so a
+        //  kernel that reads or writes past the last pixel faults.
+        ImageRGB32 file_image(m_image);
+        GuardPageImageRGB32 image(file_image);
 
         const size_t width = image.width();
         const size_t height = image.height();
@@ -303,8 +313,8 @@ class Test_ImageToBlackWhiteRGB32Range : public UnitTest{
         const uint32_t mins = uint32_t(min_color);
         const uint32_t maxs = uint32_t(max_color);
 
-        ImageRGB32 image_out(image.width(), image.height());
-        ImageRGB32 image_out_2(image.width(), image.height());
+        GuardPageImageRGB32 image_out(image.width(), image.height());
+        GuardPageImageRGB32 image_out_2(image.width(), image.height());
         size_t pixels_in_range = 0;
 
         const bool in_range_black = true;
@@ -411,7 +421,10 @@ class Test_ImageFilterByMask : public UnitTest{
     {}
 
     virtual UnitTestResult run(Logger& logger, CancellableScope& scope) const override{
-        ImageRGB32 image(m_image);
+        //  Copy the test image into guard-page memory with packed rows, so a
+        //  kernel that reads or writes past the last pixel faults.
+        ImageRGB32 file_image(m_image);
+        GuardPageImageRGB32 image(file_image);
 
         const size_t width = image.width(), height = image.height();
         cout << "Image width " << width << " height " << height << endl;
@@ -423,8 +436,8 @@ class Test_ImageFilterByMask : public UnitTest{
         compress_rgb32_to_binary_range(
             image.data(), image.bytes_per_row(), *binary_matrix, mins, maxs
         );
-        ImageRGB32 new_image = image.copy();
-        ImageRGB32 new_image_2 = image.copy();
+        GuardPageImageRGB32 new_image(file_image);
+        GuardPageImageRGB32 new_image_2(file_image);
 
         Color replacement_color = COLOR_WHITE;
         bool replace_zero_bits = true;
@@ -507,7 +520,10 @@ class Test_CompressRGB32ToBinaryEuclidean : public UnitTest{
     {}
 
     virtual UnitTestResult run(Logger& logger, CancellableScope& scope) const override{
-        ImageRGB32 image(m_image);
+        //  Copy the test image into guard-page memory with packed rows, so a
+        //  kernel that reads or writes past the last pixel faults.
+        ImageRGB32 file_image(m_image);
+        GuardPageImageRGB32 image(file_image);
 
         const size_t width = image.width();
         const size_t height = image.height();
diff --git a/SerialPrograms/Source/Kernels/ImageStats/Kernels_ImagePixelSumSqrDev_TestInfra.cpp b/SerialPrograms/Source/Kernels/ImageStats/Kernels_ImagePixelSumSqrDev_TestInfra.cpp
index d97cf0530b..bcab09234e 100644
--- a/SerialPrograms/Source/Kernels/ImageStats/Kernels_ImagePixelSumSqrDev_TestInfra.cpp
+++ b/SerialPrograms/Source/Kernels/ImageStats/Kernels_ImagePixelSumSqrDev_TestInfra.cpp
@@ -5,7 +5,7 @@
  */
 
 #include <random>
-#include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Kernels_ImagePixelSumSqrDev_TestInfra.h"
 
 namespace PokemonAutomation{
@@ -44,8 +44,10 @@ Test_SumSqrDeviation::Test_SumSqrDeviation(
 {}
 
 UnitTestResult Test_SumSqrDeviation::run(Logger& logger, CancellableScope& scope) const{
-    AlignedVector<uint32_t> ref(m_width * m_height);
-    AlignedVector<uint32_t> img(m_width * m_height);
+    //  Both images end right before a guard page, so a kernel that reads past
+    //  the last row faults on every run.
+    GuardPageBuffer<uint32_t> ref(m_width * m_height);
+    GuardPageBuffer<uint32_t> img(m_width * m_height);
 
     std::mt19937 generator(42);
     std::uniform_int_distribution<uint32_t> distribution;
diff --git a/SerialPrograms/Source/Kernels/ImageStats/Kernels_ImagePixelSumSqr_TestInfra.cpp b/SerialPrograms/Source/Kernels/ImageStats/Kernels_ImagePixelSumSqr_TestInfra.cpp
index 6e772dbeb9..f2c7c63e77 100644
--- a/SerialPrograms/Source/Kernels/ImageStats/Kernels_ImagePixelSumSqr_TestInfra.cpp
+++ b/SerialPrograms/Source/Kernels/ImageStats/Kernels_ImagePixelSumSqr_TestInfra.cpp
@@ -5,7 +5,7 @@
  */
 
 #include <random>
-#include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Kernels_ImagePixelSumSqr_TestInfra.h"
 
 namespace PokemonAutomation{
@@ -28,8 +28,10 @@ Test_PixelSumSqr::Test_PixelSumSqr(
 {}
 
 UnitTestResult Test_PixelSumSqr::run(Logger& logger, CancellableScope& scope) const{
-    AlignedVector<uint32_t> image(m_width * m_height);
-    AlignedVector<uint32_t> alpha(m_width * m_height);
+    //  Both images end right before a guard page, so a kernel that reads past
+    //  the last row faults on every run.
+    GuardPageBuffer<uint32_t> image(m_width * m_height);
+    GuardPageBuffer<uint32_t> alpha(m_width * m_height);
 
     std::mt19937 generator(42);
     std::uniform_int_distribution<uint32_t> distribution;
diff --git a/SerialPrograms/Source/Kernels/ScaleInvariantMatrixMatch/Kernels_ScaleInvariantMatrixMatch_TestInfra.cpp b/SerialPrograms/Source/Kernels/ScaleInvariantMatrixMatch/Kernels_ScaleInvariantMatrixMatch_TestInfra.cpp
index 463141fd88..e08989f3c0 100644
--- a/SerialPrograms/Source/Kernels/ScaleInvariantMatrixMatch/Kernels_ScaleInvariantMatrixMatch_TestInfra.cpp
+++ b/SerialPrograms/Source/Kernels/ScaleInvariantMatrixMatch/Kernels_ScaleInvariantMatrixMatch_TestInfra.cpp
@@ -11,6 +11,7 @@
 #include "Common/Compiler.h"
 #include "Common/Cpp/CancellableScope.h"
 #include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Kernels_ScaleInvariantMatrixMatch_TestInfra.h"
 
 namespace PokemonAutomation{
@@ -27,14 +28,9 @@ namespace ScaleInvariantMatrixMatch{
 //
 static constexpr size_t HEAD_PAD = 16;
 //
-//  Tail padding after the last row. The SSE4.1 partial-tail load
-//  (PSHUFB-based `load_f32_no_read_past_end`) always reads a full 16-byte
-//  word, which extends up to VL - 1 floats past the end of the row even
-//  though only the first `length` lanes are used. 15 floats (VL - 1 for
-//  the widest supported ISA) keeps those reads inside the allocation for
-//  the last row.
+//  There is no tail padding: the last row ends right before a guard page,
+//  so a read past the end of the matrix faults on every run.
 //
-static constexpr size_t TAIL_PAD = 15;
 
 //
 //  A result passes if it is within 4096 ULPs of the Default result, or
@@ -103,17 +99,20 @@ UnitTestResult Test::run(Logger& logger, CancellableScope& scope) const{
     //  [0, 1]. TW = T * W is precomputed with the same float rounding
     //  that production uses.
     //
-    AlignedVector<float> A(n + HEAD_PAD + TAIL_PAD);
-    AlignedVector<float> T(n + HEAD_PAD + TAIL_PAD);
+    //  Each matrix ends right before a guard page. The kernels align their
+    //  loads to T and step A and W by the same offset, so every matrix must
+    //  have the same alignment. Equal sizes give them that.
+    //
+    const size_t weighted_size = m_use_weights ? n + HEAD_PAD : 0;
+    GuardPageBuffer<float> A(n + HEAD_PAD);
+    GuardPageBuffer<float> T(n + HEAD_PAD);
     AlignedVector<const float*> rowsA(m_height);
     AlignedVector<const float*> rowsT(m_height);
-    AlignedVector<float> W;
-    AlignedVector<float> TW;
+    GuardPageBuffer<float> W(weighted_size);
+    GuardPageBuffer<float> TW(weighted_size);
     AlignedVector<const float*> rowsW;
     AlignedVector<const float*> rowsTW;
     if (m_use_weights){
-        W = AlignedVector<float>(n + HEAD_PAD + TAIL_PAD);
-        TW = AlignedVector<float>(n + HEAD_PAD + TAIL_PAD);
         rowsW = AlignedVector<const float*>(m_height);
         rowsTW = AlignedVector<const float*>(m_height);
     }
diff --git a/SerialPrograms/Source/Kernels/SpikeConvolution/Kernels_SpikeConvolution_TestInfra.cpp b/SerialPrograms/Source/Kernels/SpikeConvolution/Kernels_SpikeConvolution_TestInfra.cpp
index 11120c1d87..05c34f2f25 100644
--- a/SerialPrograms/Source/Kernels/SpikeConvolution/Kernels_SpikeConvolution_TestInfra.cpp
+++ b/SerialPrograms/Source/Kernels/SpikeConvolution/Kernels_SpikeConvolution_TestInfra.cpp
@@ -4,11 +4,13 @@
  *
  */
 
+#include <algorithm>
 #include <cstdio>
+#include <limits>
 #include <random>
 #include "Common/Compiler.h"
 #include "Common/Cpp/CancellableScope.h"
-#include "Common/Cpp/Containers/AlignedVector.tpp"
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
 #include "Kernels_SpikeConvolution_TestInfra.h"
 
 namespace PokemonAutomation{
@@ -36,11 +38,11 @@ static constexpr int MAX_WEIGHT = 8;
 static constexpr size_t MAX_LENGTH_I = 1000;
 
 //
-//  Extra room in the output buffers: the kernel may write up to
-//  PA_ALIGNMENT/sizeof(float) - 1 = 15 floats past the valid outputs
-//  (padding of the final vector).
+//  The documented size of "out": the valid outputs rounded up to 16 floats.
 //
-static constexpr size_t OUTPUT_PAD = 16;
+static size_t output_capacity(size_t lengthO){
+    return (lengthO + 15) & ~(size_t)15;
+}
 
 
 
@@ -60,19 +62,36 @@ Test::Test(
 {}
 
 UnitTestResult Test::run(Logger& logger, CancellableScope& scope) const{
+    //
+    //  "in", "kernel" and both outputs each end right before a guard page.
+    //  A read past the end of "in" or "kernel", or a write past the
+    //  documented end of "out", faults on every run. The buffers are mapped
+    //  once for the longest input. Each input length puts its data at the
+    //  end of them.
+    //
+    const size_t max_capacityO = output_capacity(MAX_LENGTH_I - m_lengthK + 1);
+    GuardPageBuffer<float> in_buffer(MAX_LENGTH_I);
+    GuardPageBuffer<float> kernel(m_lengthK);
+    GuardPageBuffer<float> out0_buffer(max_capacityO, PA_ALIGNMENT);
+    GuardPageBuffer<float> out1_buffer(max_capacityO, PA_ALIGNMENT);
+
     for (size_t lengthI = m_lengthK; lengthI <= MAX_LENGTH_I; lengthI++){
         scope.throw_if_cancelled();
         const size_t lengthO = lengthI - m_lengthK + 1;
+        const size_t capacityO = output_capacity(lengthO);
 
         //
         //  Fresh random data for each (K, N), with a deterministic seed.
-        //  "in" holds exactly lengthI floats, so any out-of-bounds read
-        //  by an implementation shows up here.
+        //  "in" holds exactly lengthI floats. "out0" and "out1" each hold
+        //  exactly capacityO floats and start 64-byte aligned (both
+        //  capacities are multiples of 16 floats). They start as NaN, so an
+        //  output that an implementation never writes shows up as a mismatch.
         //
-        AlignedVector<float> in(lengthI);
-        AlignedVector<float> kernel(m_lengthK);
-        AlignedVector<float> out0(lengthO + OUTPUT_PAD);
-        AlignedVector<float> out1(lengthO + OUTPUT_PAD);
+        float* in = in_buffer.data() + MAX_LENGTH_I - lengthI;
+        float* out0 = out0_buffer.data() + max_capacityO - capacityO;
+        float* out1 = out1_buffer.data() + max_capacityO - capacityO;
+        std::fill(out0, out0 + capacityO, std::numeric_limits<float>::quiet_NaN());
+        std::fill(out1, out1 + capacityO, std::numeric_limits<float>::quiet_NaN());
 
         std::mt19937 generator((unsigned int)(lengthI * 1000003u + m_lengthK * 7919u));
         std::uniform_int_distribution<int> data_distribution(-MAX_DATA, MAX_DATA);
@@ -84,8 +103,8 @@ UnitTestResult Test::run(Logger& logger, CancellableScope& scope) const{
             kernel[c] = (float)weight_distribution(generator);
         }
 
-        m_conv0(out0.data(), in.data(), lengthI, kernel.data(), m_lengthK);
-        m_conv1(out1.data(), in.data(), lengthI, kernel.data(), m_lengthK);
+        m_conv0(out0, in, lengthI, kernel.data(), m_lengthK);
+        m_conv1(out1, in, lengthI, kernel.data(), m_lengthK);
 
         //
         //  Both implementations compute identical integer sums, so the
diff --git a/SerialPrograms/Source/Kernels/Waterfill/Kernels_Waterfill_Tests.cpp b/SerialPrograms/Source/Kernels/Waterfill/Kernels_Waterfill_Tests.cpp
index 984a6f8939..d3ca8b7a66 100644
--- a/SerialPrograms/Source/Kernels/Waterfill/Kernels_Waterfill_Tests.cpp
+++ b/SerialPrograms/Source/Kernels/Waterfill/Kernels_Waterfill_Tests.cpp
@@ -24,6 +24,7 @@
 #include "Kernels/Waterfill/Kernels_Waterfill_Core_64xH_Default.h"
 #include "Kernels/Waterfill/Kernels_Waterfill_Routines.h"
 #include "Kernels_Waterfill_Tests.h"
+#include "Tests/GuardPageImage.h"
 #include "Tests/TestUtils.h"
 
 #include <iostream>
@@ -44,7 +45,10 @@ class Test_Waterfill : public UnitTest{
     {}
 
     virtual UnitTestResult run(Logger& logger, CancellableScope& scope) const override{
-        ImageRGB32 image(m_image);
+        //  Copy the test image into guard-page memory with packed rows, so a
+        //  kernel that reads or writes past the last pixel faults.
+        ImageRGB32 file_image(m_image);
+        GuardPageImageRGB32 image(file_image);
 
         const size_t width = image.width();
         const size_t height = image.height();
diff --git a/SerialPrograms/Source/Tests/GuardPageImage.h b/SerialPrograms/Source/Tests/GuardPageImage.h
new file mode 100644
index 0000000000..279322a0c8
--- /dev/null
+++ b/SerialPrograms/Source/Tests/GuardPageImage.h
@@ -0,0 +1,66 @@
+/*  Guard Page Image
+ *
+ *  From: https://github.com/PokemonAutomation/
+ *
+ *  An RGB32 image for unit tests of image kernels. The rows are packed with
+ *  no padding (bytes_per_row = width * 4), and the last row ends right before
+ *  a no-access guard page. A kernel that reads or writes past the last pixel
+ *  faults right away.
+ *
+ *  ImageRGB32 pads every row to 64 bytes, so tests on it never notice a kernel
+ *  that reads a little past the end of a row. Views of other memory (like a
+ *  QImage frame) don't have that padding.
+ *
+ */
+
+#ifndef PokemonAutomation_Tests_GuardPageImage_H
+#define PokemonAutomation_Tests_GuardPageImage_H
+
+#include <stdint.h>
+#include "Common/Cpp/Containers/GuardPageBuffer.h"
+#include "CommonFramework/ImageTypes/ImageViewRGB32.h"
+
+namespace PokemonAutomation{
+
+
+class GuardPageImageRGB32{
+public:
+    //  A blank image. Every pixel starts as zero.
+    GuardPageImageRGB32(size_t width, size_t height)
+        : m_width(width)
+        , m_height(height)
+        , m_pixels(width * height)
+    {}
+
+    //  A copy of "image".
+    explicit GuardPageImageRGB32(const ImageViewRGB32& image)
+        : GuardPageImageRGB32(image.width(), image.height())
+    {
+        for (size_t y = 0; y < m_height; y++){
+            for (size_t x = 0; x < m_width; x++){
+                m_pixels[x + y * m_width] = image.pixel(x, y);
+            }
+        }
+    }
+
+    size_t width() const{ return m_width; }
+    size_t height() const{ return m_height; }
+    size_t bytes_per_row() const{ return m_width * sizeof(uint32_t); }
+
+    const uint32_t* data() const{ return m_pixels.data(); }
+          uint32_t* data()      { return m_pixels.data(); }
+
+    uint32_t pixel(size_t x, size_t y) const{
+        return m_pixels[x + y * m_width];
+    }
+
+private:
+    size_t m_width;
+    size_t m_height;
+    GuardPageBuffer<uint32_t> m_pixels;
+};
+
+
+
+}
+#endif
diff --git a/SerialPrograms/cmake/SourceFiles.cmake b/SerialPrograms/cmake/SourceFiles.cmake
index d6dce27806..e91f4296a1 100644
--- a/SerialPrograms/cmake/SourceFiles.cmake
+++ b/SerialPrograms/cmake/SourceFiles.cmake
@@ -53,6 +53,8 @@ file(GLOB CORE_LIBRARY_SOURCES
     ../Common/Cpp/Containers/CircularBuffer.h
     ../Common/Cpp/Containers/FixedLimitVector.h
     ../Common/Cpp/Containers/FixedLimitVector.tpp
+    ../Common/Cpp/Containers/GuardPageBuffer.cpp
+    ../Common/Cpp/Containers/GuardPageBuffer.h
     ../Common/Cpp/Containers/Pimpl.h
     ../Common/Cpp/Containers/Pimpl.tpp
     ../Common/Cpp/CpuId/CpuId.cpp
@@ -3087,6 +3089,7 @@ file(GLOB LIBRARY_SOURCES
     Source/StaticRegistrationQt.cpp
     Source/Tests/CommandLineTests.cpp
     Source/Tests/CommandLineTests.h
+    Source/Tests/GuardPageImage.h
     Source/Tests/TestUtils.cpp
     Source/Tests/TestUtils.h
     Source/VideoGames/ZeldaTotK/Programs/ZeldaTotK_BowItemDuper.cpp