From d14e32b19aef2cfeb13a5e4b349d5855bd5e3e3a Mon Sep 17 00:00:00 2001 From: Filip Leonarski Date: Sun, 11 Oct 2026 13:44:30 +0200 Subject: [PATCH] decode kernel: __launch_bounds__(256, 5), so sm_120 keeps the five blocks its shared memory allows The compiler's register choice held the fused decode/un-transpose/preprocess kernel to four resident blocks per SM on the RTX 5080, where its shared buffer allows five. Asking for five (48 registers, a small spill) takes it from ~391 to ~313 us per 18 Mpx frame (myob, -N 1, nsys, repeated). sm_75 cannot hold five 256-thread blocks, so its code is unchanged (55 registers). Output identical: ci tier base vs new, compare 0 changed, p.hkl identical on 26 sets, MTZ reflection arrays identical on 49 files; [ImagePreprocessorGPU] passes. Co-Authored-By: Claude Opus 5.5 (1M context) Claude-Session: https://claude.ai/code/session_01EizBKhTqrAYago9KJG3dA3 --- .../image_preprocessing/ImagePreprocessorGPU.cu | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu index 1610a8658..8b4167efc 100644 --- a/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu +++ b/image_analysis/image_preprocessing/ImagePreprocessorGPU.cu @@ -262,8 +262,14 @@ __global__ __launch_bounds__(256) void untranspose_preprocess_kernel( // 16 kB blocks 32-bit detector data comes in, against the thirty-two warps the standalone kernel // keeps in flight. So this trades decode parallelism away for decode latency, and which way that // comes out is a measurement rather than an argument. +// +// Five resident blocks of 256 threads is what the shared buffer allows on an sm_120 card (100 kB per +// SM), but the compiler's register choice held it to four; asking for five takes the kernel from +// ~390 to ~315 us per 18 Mpx frame there. Where five cannot be resident anyway (sm_75: 1024 threads +// per SM) the bound changes nothing. The block is never wider than 256 threads: the shared buffer is +// at most FUSED_MAX_BLOCK_BYTES, and the launch gives it one thread per 64 bytes. template -__global__ void decode_untranspose_preprocess_kernel( +__global__ __launch_bounds__(256, 5) void decode_untranspose_preprocess_kernel( const uint8_t *__restrict__ src, const BSLZ4BlockDesc *__restrict__ desc, uint32_t *__restrict__ status,