Merge branch 'decode-bounds' into rc175: decode kernel launch bounds allow five resident blocks on sm_120

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01EizBKhTqrAYago9KJG3dA3

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<class T, int ES>
-__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,