diff --git a/CHANGELOG b/CHANGELOG
--- a/CHANGELOG
+++ b/CHANGELOG
@@ -1,5 +1,11 @@
 # Changelog
 
+- 0.3.4 (2026-11-14)
+  * The native kernels zero the stack region they use before
+    returning, so that no copies of the key remain there.
+  * On aarch64, 'cipher' and 'keystream' process inputs of 256 bytes
+    or more four blocks at a time, roughly twice as fast as before.
+
 - 0.3.3 (2026-11-14)
   * Adds four-block kernels producing interleaved keystream blocks,
     as used by ChaCha8Rand, and roughly twice as fast as the
diff --git a/cbits/chacha20_native.c b/cbits/chacha20_native.c
--- a/cbits/chacha20_native.c
+++ b/cbits/chacha20_native.c
@@ -19,8 +19,45 @@
     0x61707865u, 0x3320646eu, 0x79622d32u, 0x6b206574u
 };
 
+/*
+ * A kernel's stack frame holds copies of the key, e.g. in spilled
+ * state registers, after it returns.  Each exported entry point thus
+ * runs its kernel in a frame of its own and then overwrites the stack
+ * region below it with zeros, so that no key material remains there.
+ *
+ * The region must exceed the kernels' frames, including any they call,
+ * with room for other compilers and flags.  At -O2 (clang 21, gcc 14
+ * and 16), the single- and two-block kernels' frames are at most
+ * about 400 bytes, and the four- and eight-block kernels' at most
+ * about 700 bytes on aarch64 and 2.2 KiB on x86_64.
+ */
+#define BURN_SMALL 1024
+#if defined(__aarch64__)
+#define BURN_LARGE 2048
+#else
+#define BURN_LARGE 4096
 #endif
 
+/* zero 'n' bytes, in a way the compiler may not elide */
+static inline void scrub(uint8_t *buf, size_t n) {
+    memset(buf, 0, n);
+    __asm__ __volatile__("" : : "r"(buf) : "memory");
+}
+
+__attribute__((noinline))
+static void burn_stack_small(void) {
+    uint8_t buf[BURN_SMALL];
+    scrub(buf, sizeof buf);
+}
+
+__attribute__((noinline))
+static void burn_stack_large(void) {
+    uint8_t buf[BURN_LARGE];
+    scrub(buf, sizeof buf);
+}
+
+#endif
+
 #if defined(__aarch64__)
 
 /*
@@ -91,6 +128,42 @@
     *v3 = vaddq_u32(d, s3);
 }
 
+/*
+ * ChaCha NEON kernel using inter-block parallelism: four consecutive
+ * blocks at once, with register x[i] holding word i of each block.  A
+ * quarter-round is then one set of element-wise operations on four
+ * registers, with no lane rotation, and the four quarter-rounds of
+ * each half-round are independent chains.  The registers, stored in
+ * order, give the blocks interleaved word by word.
+ */
+static inline __attribute__((always_inline))
+void chacha20_core4v(uint32x4_t x[16], int drounds) {
+    for (int i = 0; i < drounds; i++) {
+        QUARTER(x[0], x[4], x[ 8], x[12]);
+        QUARTER(x[1], x[5], x[ 9], x[13]);
+        QUARTER(x[2], x[6], x[10], x[14]);
+        QUARTER(x[3], x[7], x[11], x[15]);
+        QUARTER(x[0], x[5], x[10], x[15]);
+        QUARTER(x[1], x[6], x[11], x[12]);
+        QUARTER(x[2], x[7], x[ 8], x[13]);
+        QUARTER(x[3], x[4], x[ 9], x[14]);
+    }
+}
+
+/* Transpose four registers, each holding one word of four blocks,
+ * into four each holding four consecutive words of one block. */
+static inline void chacha20_transpose4(uint32x4_t *a, uint32x4_t *b,
+                                       uint32x4_t *c, uint32x4_t *d) {
+    uint64x2_t t0 = vreinterpretq_u64_u32(vtrn1q_u32(*a, *b));
+    uint64x2_t t1 = vreinterpretq_u64_u32(vtrn2q_u32(*a, *b));
+    uint64x2_t t2 = vreinterpretq_u64_u32(vtrn1q_u32(*c, *d));
+    uint64x2_t t3 = vreinterpretq_u64_u32(vtrn2q_u32(*c, *d));
+    *a = vreinterpretq_u32_u64(vtrn1q_u64(t0, t2));
+    *b = vreinterpretq_u32_u64(vtrn1q_u64(t1, t3));
+    *c = vreinterpretq_u32_u64(vtrn2q_u64(t0, t2));
+    *d = vreinterpretq_u32_u64(vtrn2q_u64(t1, t3));
+}
+
 /* Set up the constant rows of the state from key + nonce.  s3
  * (counter + nonce) varies per block and is built inside the loop. */
 static inline void chacha20_setup(const uint8_t key[32],
@@ -111,7 +184,8 @@
  * Generate one 64-byte keystream block at 'out', using 'drounds'
  * double rounds (10 for ChaCha20).
  */
-void chacha20_block_native(const uint8_t key[32], uint32_t counter,
+static __attribute__((noinline))
+void chacha20_block_kernel(const uint8_t key[32], uint32_t counter,
                            const uint8_t nonce[12], uint8_t out[64],
                            int drounds) {
     uint32x4_t s0, s1, s2;
@@ -130,12 +204,58 @@
 }
 
 /*
+ * Encrypt/decrypt the longest prefix of 'in' that is a whole number of
+ * 256-byte groups of four blocks, via the four-block core, transposing
+ * each group's rows back into block order.  Returns the number of
+ * bytes processed.  'in' and 'out' may be the same buffer.
+ */
+static __attribute__((noinline))
+size_t chacha20_cipher4_kernel(const uint8_t key[32], uint32_t counter,
+                               const uint8_t nonce[12],
+                               const uint8_t *in, uint8_t *out,
+                               size_t inlen, int drounds) {
+    uint32_t k[8], n[3];
+    memcpy(k, key, 32);
+    memcpy(n, nonce, 12);
+    uint32x4_t s[16];
+    for (int i = 0; i < 4; i++) s[i] = vdupq_n_u32(chacha_constants[i]);
+    for (int i = 0; i < 8; i++) s[4 + i] = vdupq_n_u32(k[i]);
+    for (int i = 0; i < 3; i++) s[13 + i] = vdupq_n_u32(n[i]);
+    static const uint32_t lanes[4] = { 0, 1, 2, 3 };
+    const uint32x4_t inc = vld1q_u32(lanes);
+
+    size_t pos = 0;
+    while (pos + 256 <= inlen) {
+        s[12] = vaddq_u32(vdupq_n_u32(counter), inc);
+        uint32x4_t x[16];
+        for (int i = 0; i < 16; i++) x[i] = s[i];
+        chacha20_core4v(x, drounds);
+        for (int i = 0; i < 16; i++) x[i] = vaddq_u32(x[i], s[i]);
+
+        for (int r = 0; r < 4; r++) {
+            chacha20_transpose4(&x[4 * r], &x[4 * r + 1],
+                                &x[4 * r + 2], &x[4 * r + 3]);
+            for (int j = 0; j < 4; j++) {
+                size_t o = pos + 64 * j + 16 * r;
+                vst1q_u8(out + o,
+                         veorq_u8(vld1q_u8(in + o),
+                                  vreinterpretq_u8_u32(x[4 * r + j])));
+            }
+        }
+        pos += 256;
+        counter += 4;
+    }
+    return pos;
+}
+
+/*
  * Encrypt/decrypt 'inlen' bytes at 'in' into 'out' using ChaCha with
  * 'drounds' double rounds (10 for ChaCha20) and the given key,
  * starting counter, and nonce.  Stream cipher, so the same routine
  * decrypts.  'in' and 'out' may be the same buffer.
  */
-void chacha20_cipher_native(const uint8_t key[32], uint32_t counter,
+static __attribute__((noinline))
+void chacha20_cipher_kernel(const uint8_t key[32], uint32_t counter,
                             const uint8_t nonce[12],
                             const uint8_t *in, uint8_t *out,
                             size_t inlen, int drounds) {
@@ -189,27 +309,6 @@
 }
 
 /*
- * ChaCha NEON kernel using inter-block parallelism: four consecutive
- * blocks at once, with register x[i] holding word i of each block.  A
- * quarter-round is then one set of element-wise operations on four
- * registers, with no lane rotation, and the four quarter-rounds of
- * each half-round are independent chains.  The registers, stored in
- * order, give the blocks interleaved word by word.
- */
-static inline void chacha20_core4v(uint32x4_t x[16], int drounds) {
-    for (int i = 0; i < drounds; i++) {
-        QUARTER(x[0], x[4], x[ 8], x[12]);
-        QUARTER(x[1], x[5], x[ 9], x[13]);
-        QUARTER(x[2], x[6], x[10], x[14]);
-        QUARTER(x[3], x[7], x[11], x[15]);
-        QUARTER(x[0], x[5], x[10], x[15]);
-        QUARTER(x[1], x[6], x[11], x[12]);
-        QUARTER(x[2], x[7], x[ 8], x[13]);
-        QUARTER(x[3], x[4], x[ 9], x[14]);
-    }
-}
-
-/*
  * Write 'ngroups' groups of four consecutive keystream blocks to
  * 'out', 256 bytes per group, using 'drounds' double rounds and
  * starting at block 'counter'.  Within a group the blocks are
@@ -219,7 +318,8 @@
  * public constant, counter, and nonce words, as in ChaCha8Rand.
  * 'out' must not overlap 'key' or 'nonce'.
  */
-void chacha20_interleaved_native(const uint8_t key[32], uint32_t counter,
+static __attribute__((noinline))
+void chacha20_interleaved_kernel(const uint8_t key[32], uint32_t counter,
                                  const uint8_t nonce[12], uint8_t *out,
                                  size_t ngroups, int drounds,
                                  int key_ff_only) {
@@ -250,6 +350,42 @@
     }
 }
 
+/* exported entry points: run each kernel, then scrub the stack */
+
+void chacha20_block_native(const uint8_t key[32], uint32_t counter,
+                           const uint8_t nonce[12], uint8_t out[64],
+                           int drounds) {
+    chacha20_block_kernel(key, counter, nonce, out, drounds);
+    burn_stack_small();
+}
+
+void chacha20_cipher_native(const uint8_t key[32], uint32_t counter,
+                            const uint8_t nonce[12],
+                            const uint8_t *in, uint8_t *out,
+                            size_t inlen, int drounds) {
+    /* groups of four blocks via the four-block kernel, whose frame is
+     * deeper, and the rest via the single-block kernel */
+    size_t done = 0;
+    if (inlen >= 256)
+        done = chacha20_cipher4_kernel(key, counter, nonce, in, out,
+                                       inlen, drounds);
+    chacha20_cipher_kernel(key, counter + (uint32_t)(done / 64), nonce,
+                           in + done, out + done, inlen - done, drounds);
+    if (done)
+        burn_stack_large();
+    else
+        burn_stack_small();
+}
+
+void chacha20_interleaved_native(const uint8_t key[32], uint32_t counter,
+                                 const uint8_t nonce[12], uint8_t *out,
+                                 size_t ngroups, int drounds,
+                                 int key_ff_only) {
+    chacha20_interleaved_kernel(key, counter, nonce, out, ngroups,
+                                drounds, key_ff_only);
+    burn_stack_large();
+}
+
 int chacha20_native_available(void) {
     return 1;
 }
@@ -558,7 +694,8 @@
  * Generate one 64-byte keystream block at 'out', using 'drounds'
  * double rounds (10 for ChaCha20).
  */
-void chacha20_block_native(const uint8_t key[32], uint32_t counter,
+static __attribute__((noinline))
+void chacha20_block_kernel(const uint8_t key[32], uint32_t counter,
                            const uint8_t nonce[12], uint8_t out[64],
                            int drounds) {
     __m128i s0, s1, s2;
@@ -581,7 +718,8 @@
  * starting counter, and nonce.  Stream cipher, so the same routine
  * decrypts.  'in' and 'out' may be the same buffer.
  */
-void chacha20_cipher_native(const uint8_t key[32], uint32_t counter,
+static __attribute__((noinline))
+void chacha20_cipher_kernel(const uint8_t key[32], uint32_t counter,
                             const uint8_t nonce[12],
                             const uint8_t *in, uint8_t *out,
                             size_t inlen, int drounds) {
@@ -641,7 +779,7 @@
 
 /* ChaCha core over eight consecutive blocks, one state word per
  * register, as in the four-block SSE2 kernel on wider registers. */
-__attribute__((target("avx2")))
+__attribute__((target("avx2"), always_inline))
 static inline void chacha20_core8v_avx2(__m256i x[16], int drounds) {
     for (int i = 0; i < drounds; i++) {
         QUARTER_AVX2(x[0], x[4], x[ 8], x[12]);
@@ -661,7 +799,7 @@
  * register holds the first group's word, the high lane the second's.
  * Returns the number of groups written.
  */
-__attribute__((target("avx2")))
+__attribute__((target("avx2"), noinline))
 static size_t chacha20_interleaved_avx2(const uint8_t key[32],
                                         uint32_t counter,
                                         const uint8_t nonce[12],
@@ -704,7 +842,8 @@
  * kernel of the same name: register x[i] holds word i of each of
  * four consecutive blocks.
  */
-static inline void chacha20_core4v(__m128i x[16], int drounds) {
+static inline __attribute__((always_inline))
+void chacha20_core4v(__m128i x[16], int drounds) {
     for (int i = 0; i < drounds; i++) {
         QUARTER(x[0], x[4], x[ 8], x[12]);
         QUARTER(x[1], x[5], x[ 9], x[13]);
@@ -727,7 +866,8 @@
  * public constant, counter, and nonce words, as in ChaCha8Rand.
  * 'out' must not overlap 'key' or 'nonce'.
  */
-void chacha20_interleaved_native(const uint8_t key[32], uint32_t counter,
+static __attribute__((noinline))
+void chacha20_interleaved_kernel(const uint8_t key[32], uint32_t counter,
                                  const uint8_t nonce[12], uint8_t *out,
                                  size_t ngroups, int drounds,
                                  int key_ff_only) {
@@ -735,16 +875,6 @@
     memcpy(k, key, 32);
     memcpy(n, nonce, 12);
 
-    /* pairs of groups via AVX2 where available, the rest via SSE2 */
-    if (ngroups >= 2 && have_avx2()) {
-        size_t done = chacha20_interleaved_avx2(key, counter, nonce, out,
-                                                ngroups, drounds,
-                                                key_ff_only);
-        counter += (uint32_t)(4 * done);
-        out     += 256 * done;
-        ngroups -= done;
-    }
-
     __m128i s[16];
     for (int i = 0; i < 4; i++)
         s[i] = _mm_set1_epi32((int)chacha_constants[i]);
@@ -766,6 +896,39 @@
         }
         counter += 4;
     }
+}
+
+/* exported entry points: run each kernel, then scrub the stack */
+
+void chacha20_block_native(const uint8_t key[32], uint32_t counter,
+                           const uint8_t nonce[12], uint8_t out[64],
+                           int drounds) {
+    chacha20_block_kernel(key, counter, nonce, out, drounds);
+    burn_stack_small();
+}
+
+void chacha20_cipher_native(const uint8_t key[32], uint32_t counter,
+                            const uint8_t nonce[12],
+                            const uint8_t *in, uint8_t *out,
+                            size_t inlen, int drounds) {
+    chacha20_cipher_kernel(key, counter, nonce, in, out, inlen, drounds);
+    burn_stack_small();
+}
+
+void chacha20_interleaved_native(const uint8_t key[32], uint32_t counter,
+                                 const uint8_t nonce[12], uint8_t *out,
+                                 size_t ngroups, int drounds,
+                                 int key_ff_only) {
+    /* pairs of groups via AVX2 where available, the rest via SSE2, each
+     * kernel called from here so that their frames don't stack */
+    size_t done = 0;
+    if (ngroups >= 2 && have_avx2())
+        done = chacha20_interleaved_avx2(key, counter, nonce, out,
+                                         ngroups, drounds, key_ff_only);
+    chacha20_interleaved_kernel(key, counter + (uint32_t)(4 * done),
+                                nonce, out + 256 * done, ngroups - done,
+                                drounds, key_ff_only);
+    burn_stack_large();
 }
 
 int chacha20_native_available(void) {
diff --git a/ppad-chacha.cabal b/ppad-chacha.cabal
--- a/ppad-chacha.cabal
+++ b/ppad-chacha.cabal
@@ -1,6 +1,6 @@
 cabal-version:      3.0
 name:               ppad-chacha
-version:            0.3.3
+version:            0.3.4
 synopsis:           A fast ChaCha20 stream cipher
 license:            MIT
 license-file:       LICENSE
@@ -57,6 +57,7 @@
   hs-source-dirs:      test
   main-is:             Main.hs
   other-modules:       Vectors
+  c-sources:           test/remnant.c
 
   ghc-options:
     -rtsopts -Wall -O2
diff --git a/test/Main.hs b/test/Main.hs
--- a/test/Main.hs
+++ b/test/Main.hs
@@ -18,6 +18,7 @@
 import Data.Maybe (fromJust)
 import qualified Data.Primitive.PrimArray as PA
 import Data.Word (Word8, Word32)
+import Foreign.C.Types (CInt(..))
 import Foreign.Ptr (Ptr, castPtr)
 import Test.Tasty
 import qualified Test.Tasty.HUnit as H
@@ -47,6 +48,7 @@
       , interleaved_layout
       , interleaved_dispatch
       ]
+  , stack_scrub
   ]
 
 quarter :: TestTree
@@ -344,3 +346,28 @@
     ref <- interleave Reduced._interleaved_pure_into r ff key con
              block_non groups
     H.assertEqual (show (r, ff, con, groups)) ref out
+
+-- stack scrubbing
+
+foreign import ccall unsafe "chacha_test_remnant"
+  c_test_remnant :: CInt -> CInt -> IO CInt
+
+-- no native kernel leaves its key in the stack region it used, for
+-- each round count (see test/remnant.c); on platforms without native
+-- kernels the check is vacuous
+stack_scrub :: TestTree
+stack_scrub = H.testCase "native kernels scrub the stack" $
+  for_ [0 .. 6] $ \which ->
+  for_ [4, 6, 10] $ \drounds -> do
+    r <- c_test_remnant which drounds
+    H.assertEqual (show (kernel which, drounds)) 0 r
+  where
+    kernel :: CInt -> String
+    kernel w = case w of
+      0 -> "block"
+      1 -> "cipher (100B)"
+      2 -> "cipher (300B)"
+      3 -> "cipher (1024B)"
+      4 -> "one interleaved group"
+      5 -> "three interleaved groups"
+      _ -> "four interleaved groups"
diff --git a/test/remnant.c b/test/remnant.c
new file mode 100644
--- /dev/null
+++ b/test/remnant.c
@@ -0,0 +1,90 @@
+#include <stddef.h>
+#include <stdint.h>
+#include <string.h>
+
+void chacha20_block_native(const uint8_t *key, uint32_t counter,
+                           const uint8_t *nonce, uint8_t *out,
+                           int drounds);
+void chacha20_cipher_native(const uint8_t *key, uint32_t counter,
+                            const uint8_t *nonce, const uint8_t *in,
+                            uint8_t *out, size_t inlen, int drounds);
+void chacha20_interleaved_native(const uint8_t *key, uint32_t counter,
+                                 const uint8_t *nonce, uint8_t *out,
+                                 size_t ngroups, int drounds,
+                                 int key_ff_only);
+
+#define SCAN 8192
+
+/* the eight key words, distinct and without repeated bytes */
+static const uint32_t key_words[8] = {
+    0x1b2c3d4e, 0x5f607182, 0x93a4b5c6, 0xd7e8f90a,
+    0x2b3c4d5e, 0x6f708192, 0xa3b4c5d6, 0xe7f8091a
+};
+
+/* Does the region contain a key half (four consecutive key words, as
+ * a kernel holding a row of the state would store them) or a key word
+ * repeated four times (as one holding a broadcast word would)? */
+static int has_remnant(const volatile uint8_t *p, size_t n) {
+    for (size_t i = 0; i + 16 <= n; i++) {
+        uint32_t w[4];
+        for (int j = 0; j < 4; j++) {
+            uint8_t b[4];
+            for (int k = 0; k < 4; k++) b[k] = p[i + 4 * j + k];
+            memcpy(&w[j], b, 4);
+        }
+        for (int h = 0; h < 8; h += 4)
+            if (w[0] == key_words[h]     && w[1] == key_words[h + 1] &&
+                w[2] == key_words[h + 2] && w[3] == key_words[h + 3])
+                return 1;
+        for (int k = 0; k < 8; k++)
+            if (w[0] == key_words[k] && w[1] == key_words[k] &&
+                w[2] == key_words[k] && w[3] == key_words[k])
+                return 1;
+    }
+    return 0;
+}
+
+/* Scan the stack region below this frame, where the kernel's frame
+ * was, after calling it. */
+__attribute__((noinline))
+static int scan_below(void) {
+    volatile uint8_t here = 0;
+    const volatile uint8_t *top = &here;
+    return has_remnant(top - SCAN, SCAN);
+}
+
+static uint8_t key[32], nonce[12], buf[1024];
+
+__attribute__((noinline))
+static void call_kernel(int which, int drounds) {
+    switch (which) {
+    case 0: chacha20_block_native(key, 1, nonce, buf, drounds); break;
+    case 1: chacha20_cipher_native(key, 1, nonce, buf, buf, 100,
+                                   drounds); break;
+    case 2: chacha20_cipher_native(key, 1, nonce, buf, buf, 300,
+                                   drounds); break;
+    case 3: chacha20_cipher_native(key, 1, nonce, buf, buf, 1024,
+                                   drounds); break;
+    case 4: chacha20_interleaved_native(key, 0, nonce, buf, 1, drounds,
+                                        1); break;
+    case 5: chacha20_interleaved_native(key, 0, nonce, buf, 3, drounds,
+                                        1); break;
+    default: chacha20_interleaved_native(key, 0, nonce, buf, 4,
+                                         drounds, 1); break;
+    }
+}
+
+/*
+ * Call the given kernel (0: block; 1, 2, 3: cipher of 100, 300, and
+ * 1024 bytes; 4, 5, 6: one, three, and four interleaved groups) with a
+ * known key, and report whether any of the key remains on the stack
+ * afterwards.  call_kernel and scan_below are called from the same
+ * frame, so the scanned region covers the kernel's.
+ */
+int chacha_test_remnant(int which, int drounds) {
+    memcpy(key, key_words, 32);
+    memset(nonce, 0, sizeof nonce);
+    memset(buf, 0, sizeof buf);
+    call_kernel(which, drounds);
+    return scan_below();
+}