mx/Jungfraujoch
2
1
Fork 0
Code Issues Pull Requests Actions 4 Packages Projects Releases 67 Wiki Activity

image_preprocessing: write whole elements out of the un-transpose

Browse Source 
The raw-bytes path assembled each element a byte at a time, which on a full frame cost
about 4x against writing the 8 contiguous elements a thread owns through an
element-typed pointer. They are 8*ES-byte aligned, so the compiler merges them.
72.4 MB frame: 1.524 -> 0.406 ms for upload plus both kernels.

The test now also times the LZ4 pass on its own, so the bounds and validity checks in
the hot loop can be costed rather than guessed at. They are free: 0.231 ms against
0.2297 ms measured for the kernel before any of them existed - the restored offset == 1
and power-of-two fast paths pay for them. compute-sanitizer memcheck reports no error
over 400 single-bit-corrupted payloads and nine malformed containers.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
This commit is contained in:
leonarski_f
2026-08-03 14:16:58 +02:00
co-authored by Claude Opus 5
parent b6c4a59d69
commit 2d3c39c9dd
2 changed files with 22 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files
image_analysis/image_preprocessing/BSLZ4DecoderGPU.cu
+13 -3
View File
@@ -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;
}
}
}
tests/BSLZ4DecoderGPUFuzzTest.cpp
+9
View File
@@ -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",