diff --git a/image_analysis/image_preprocessing/BSLZ4DecoderGPU.cu b/image_analysis/image_preprocessing/BSLZ4DecoderGPU.cu
index 1e4ddab9..8068328d 100644
--- a/image_analysis/image_preprocessing/BSLZ4DecoderGPU.cu
+++ b/image_analysis/image_preprocessing/BSLZ4DecoderGPU.cu
@@ -2,6 +2,7 @@
 // SPDX-License-Identifier: GPL-3.0-only
 
 #include <algorithm>
+#include <type_traits>
 
 #include "BSLZ4DecoderGPU.h"
 #include "../../common/JFJochException.h"
@@ -163,6 +164,12 @@ namespace {
         uint8_t *dst = out + desc[b].out_off;
         const uint32_t n = size / 8;
 
+        // The 8 elements a thread owns are contiguous and 8*ES-byte aligned, so they are assembled
+        // whole and written through an element-typed pointer. Storing them byte by byte instead
+        // costs about 4x on a full frame.
+        using UT = typename std::conditional<ES == 1, uint8_t,
+                   typename std::conditional<ES == 2, uint16_t, uint32_t>::type>::type;
+
         for (uint32_t i = threadIdx.x; i < n; i += blockDim.x) {
             uint64_t x[ES];
             #pragma unroll
@@ -173,11 +180,14 @@ namespace {
                 for (int k = 0; k < 8; k++) a |= (uint64_t)pin[k * n + i] << (8 * k);
                 x[p] = transpose8(a);
             }
+            UT *dstT = reinterpret_cast<UT *>(dst) + i * 8;
             #pragma unroll
-            for (int k = 0; k < 8; k++)
+            for (int k = 0; k < 8; k++) {
+                UT v = 0;
                 #pragma unroll
-                for (int p = 0; p < ES; p++)
-                    dst[(i * 8 + k) * ES + p] = (uint8_t)(x[p] >> (8 * k));
+                for (int p = 0; p < ES; p++) v |= (UT)((UT)((x[p] >> (8 * k)) & 0xff) << (8 * p));
+                dstT[k] = v;
+            }
         }
     }
 
diff --git a/tests/BSLZ4DecoderGPUFuzzTest.cpp b/tests/BSLZ4DecoderGPUFuzzTest.cpp
index 289019e0..6fe6d097 100644
--- a/tests/BSLZ4DecoderGPUFuzzTest.cpp
+++ b/tests/BSLZ4DecoderGPUFuzzTest.cpp
@@ -678,6 +678,15 @@ TEST_CASE("BSLZ4Fuzz_LargeFrame", "[BSLZ4Fuzz]") {
     REQUIRE(cudaStreamSynchronize(*stream) == cudaSuccess);
     auto t3 = std::chrono::steady_clock::now();
 
+    // The LZ4 pass on its own (upload + parse), so the bounds and validity checks in the hot loop
+    // can be costed against the un-transpose rather than hidden behind it.
+    auto t4 = std::chrono::steady_clock::now();
+    for (int i = 0; i < reps; i++) decoder.DecodeShuffled(image);
+    REQUIRE(cudaStreamSynchronize(*stream) == cudaSuccess);
+    auto t5 = std::chrono::steady_clock::now();
+    printf("[BSLZ4Fuzz] LZ4 pass alone (upload + parse) %.3f ms\n",
+           std::chrono::duration<double, std::milli>(t5 - t4).count() / reps);
+
     const double cpu_ms = std::chrono::duration<double, std::milli>(t1 - t0).count();
     const double gpu_ms = std::chrono::duration<double, std::milli>(t3 - t2).count() / reps;
     printf("[BSLZ4Fuzz] large frame %zux%zu uint32 = %.1f MB, compressed %.1f MB (%.2fx), %zu blocks\n",