chore[gpu]: don't inline bitunpack lane impls (#7441)

Fully inlining the bitunpack kernels adds memory pressure such that the
dynamic dispatch kernel previously operated at the 32 registers per
thread max. Preventing to inline the bit unpack line impls drops the
register count per thread to 24. This is relevant for future changes
which will requires more registers such as patches support on the GPU.

The performance hit we take for bitunpack kernels here is 10%. It's
probably worthwhile to investigate whether there might a tradeoff here
to get similar perf with less aggressive inlining in the future. One
thing we could also look at is trading in register spills via launch
bounds for more occupancy.

Splitting out the lane implementations into headers is needed such that
the dynamic dispatch kernel ptx can be compiled without the standalone
bitunpack kernels which are not relevant for the dynamic dispatch
kernel. This reduces the amount of assembly for the dynamic dispatch to
48k lines from 128k lines. Besides nvcc compile times, this is relevant
for the dynamic dispatch kernel in terms of ptx to device compilation
which should be as fast as possible.

---------

Signed-off-by: Alexander Droste <alexander.droste@protonmail.com>

Author: 0ax1

.github/workflows/ci.yml

@@ -277,6 +277,7 @@ jobs:
           git ls-files vortex-cuda vortex-cxx vortex-duckdb vortex-ffi \
             | grep -E '\.(cpp|hpp|cu|cuh|h)$' \
             | grep -v 'kernels/src/bit_unpack_.*\.cu$' \
+            | grep -v 'kernels/src/bit_unpack_.*_lanes\.cuh$' \
             | xargs clang-format --dry-run --Werror --style=file
 
   rust-lint-no-default:

vortex-cuda/build.rs

@@ -14,7 +14,8 @@ use std::process::Command;
 
 use fastlanes::FastLanes;
 
-use crate::bit_unpack_gen::generate_cuda_unpack;
+use crate::bit_unpack_gen::generate_cuda_unpack_kernels;
+use crate::bit_unpack_gen::generate_cuda_unpack_lanes;
 
 #[path = "src/bit_unpack_gen.rs"]
 pub mod bit_unpack_gen;
@@ -72,7 +73,11 @@ fn main() {
 
             match path.extension().and_then(|e| e.to_str()) {
                 Some("cuh") | Some("h") => {
-                    println!("cargo:rerun-if-changed={}", path.display())
+                    // Only watch hand-written .cuh/.h files, not generated ones
+                    // (generated files are rebuilt when cuda_kernel_generator changes)
+                    if !is_generated {
+                        println!("cargo:rerun-if-changed={}", path.display());
+                    }
                 }
                 Some("cu") => {
                     // Only watch hand-written .cu files, not generated ones
@@ -94,10 +99,19 @@ fn main() {
 }
 
 fn generate_unpack<T: FastLanes>(output_dir: &Path, thread_count: usize) -> io::Result<PathBuf> {
-    let path = output_dir.join(format!("bit_unpack_{}.cu", T::T));
-    let mut cu_file = File::create(&path)?;
-    generate_cuda_unpack::<T>(&mut cu_file, thread_count)?;
-    Ok(path)
+    // Generate the lanes header (.cuh) — device functions only, no __global__ kernels.
+    // This is what dynamic_dispatch.cu includes (via bit_unpack.cuh).
+    let cuh_path = output_dir.join(format!("bit_unpack_{}_lanes.cuh", T::T));
+    let mut cuh_file = File::create(&cuh_path)?;
+    generate_cuda_unpack_lanes::<T>(&mut cuh_file)?;
+
+    // Generate the standalone kernels (.cu) — includes the lanes header,
+    // adds _device template + __global__ wrappers. Compiled to its own PTX.
+    let cu_path = output_dir.join(format!("bit_unpack_{}.cu", T::T));
+    let mut cu_file = File::create(&cu_path)?;
+    generate_cuda_unpack_kernels::<T>(&mut cu_file, thread_count)?;
+
+    Ok(cu_path)
 }
 
 fn nvcc_compile_ptx(

vortex-cuda/kernels/src/bit_unpack.cuh

@@ -7,10 +7,10 @@
 #include <cuda_runtime.h>
 #include <stdint.h>
 
-#include "bit_unpack_8.cu"
-#include "bit_unpack_16.cu"
-#include "bit_unpack_32.cu"
-#include "bit_unpack_64.cu"
+#include "bit_unpack_8_lanes.cuh"
+#include "bit_unpack_16_lanes.cuh"
+#include "bit_unpack_32_lanes.cuh"
+#include "bit_unpack_64_lanes.cuh"
 #include "patches.h"
 
 /// Decodes a single lane of packed data.
@@ -26,22 +26,22 @@
 /// * `lane` - Lane index within the block (used to determine which packed words to process)
 /// * `bit_width` - Number of bits with which each value is encoded
 template <typename T>
-__device__ inline void bit_unpack_lane(const T *__restrict packed_chunk,
-                                       T *__restrict output_buffer,
-                                       T reference,
-                                       unsigned int lane,
-                                       uint32_t bit_width);
+__device__ __noinline__ void bit_unpack_lane(const T *__restrict packed_chunk,
+                                             T *__restrict output_buffer,
+                                             T reference,
+                                             unsigned int lane,
+                                             uint32_t bit_width);
 
 /// Template specializations for `bitunpack_lane_to_smem` for different integer types.
 ///
 /// Generates template specializations for each supported integer size (8, 16, 32, 64 bits).
 #define BIT_UNPACK_LANE(bits)                                                                                \
     template <>                                                                                              \
-    __device__ inline void bit_unpack_lane<uint##bits##_t>(const uint##bits##_t *in,                         \
-                                                           uint##bits##_t *out,                              \
-                                                           uint##bits##_t reference,                         \
-                                                           unsigned int lane,                                \
-                                                           uint32_t bw) {                                    \
+    __device__ __noinline__ void bit_unpack_lane<uint##bits##_t>(const uint##bits##_t *in,                   \
+                                                                 uint##bits##_t *out,                        \
+                                                                 uint##bits##_t reference,                   \
+                                                                 unsigned int lane,                          \
+                                                                 uint32_t bw) {                              \
         bit_unpack_##bits##_lane(in, out, reference, lane, bw);                                              \
     }