Drop cub::ThreadReduce from the device sort (fixes 0.0.9 render crash)

pyproject.toml

@@ -104,9 +104,14 @@ repair-wheel-command = [
 before-all = [
   'dnf config-manager --add-repo https://developer.download.nvidia.com/compute/cuda/repos/rhel8/$(uname -m)/cuda-rhel8.repo',
   'dnf clean all',
-  'dnf install -y --setopt=install_weak_deps=False cuda-nvrtc-devel-12-8 cuda-cudart-devel-12-8 cuda-libraries-devel-12-8',
+  'dnf install -y --setopt=install_weak_deps=False cuda-nvrtc-devel-12-8 cuda-cudart-devel-12-8 cuda-libraries-devel-12-8 cuda-cccl-12-8',
   'dnf download cuda-nvcc-12-8',
   'rpm -ivh --nodeps cuda-nvcc-12-8-*.rpm',
+  # Pin /usr/local/cuda to the 12.8 toolkit just installed so the build resolves
+  # a single, consistent CUDA version -- compile flags and the CCCL/CUB headers
+  # bundled into the wheel -- instead of whatever a pre-existing /usr/local/cuda
+  # symlink happens to point at.
+  'ln -sfn /usr/local/cuda-12.8 /usr/local/cuda',
 ]
 repair-wheel-command = [
   # Do not vendor the CUDA JIT libs; they are provided by the nvidia-*-cu12

src/bridge/CMakeLists.txt

@@ -140,5 +140,17 @@ if (SKBUILD)
     endforeach()
     install(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/../data DESTINATION gs_madrona)
 
+    # Bundle the toolkit's CCCL/CUB/Thrust headers; nvrtc compiles the device
+    # code against these at first render. Report the CUB version so any skew
+    # from the device code's required API (see mw/device/sort_archetype.cpp) is
+    # visible in the build log.
+    foreach(_cuda_inc ${CUDAToolkit_INCLUDE_DIRS})
+        if (EXISTS "${_cuda_inc}/cub/version.cuh")
+            file(STRINGS "${_cuda_inc}/cub/version.cuh" _cub_version_line
+                REGEX "^#define CUB_VERSION ")
+            message(STATUS "Bundling CUDA ${CUDAToolkit_VERSION} headers (${_cub_version_line})")
+        endif()
+    endforeach()
+
     install(DIRECTORY ${CUDAToolkit_INCLUDE_DIRS} DESTINATION gs_madrona/cuda)
 endif()

src/mw/device/sort_archetype.cpp

@@ -157,31 +157,10 @@ struct BlockRadixRankMatchEarlyCountsCustom
                 atomicAdd(&warp_histograms[Digit(keys[u])][part], 1);
             }
 
-            // sum different parts;
-            // no extra work is necessary if NUM_PARTS == 1
-            if (NUM_PARTS > 1)
-            {
-                __syncwarp(WARP_MASK);
-                // TODO: handle RADIX_DIGITS % WARP_THREADS != 0 if it becomes necessary
-                const int WARP_BINS_PER_THREAD = RADIX_DIGITS / WARP_THREADS;
-                int bins[WARP_BINS_PER_THREAD];
-                #pragma unroll
-                for (int u = 0; u < WARP_BINS_PER_THREAD; ++u)
-                {
-                    int bin = lane + u * WARP_THREADS;
-                    bins[u] = ThreadReduce(warp_histograms[bin], ::cuda::std::plus<>{});
-                }
-                __syncthreads();
-
-                // store the resulting histogram in shared memory
-                int* warp_offsets = &s.warp_offsets[warp][0];
-                #pragma unroll
-                for (int u = 0; u < WARP_BINS_PER_THREAD; ++u)
-                {
-                    int bin = lane + u * WARP_THREADS;
-                    warp_offsets[bin] = bins[u];
-                }
-            }
+            // Summing across parts would go here, but every instantiation uses
+            // NUM_PARTS == 1, so there is nothing to do. (This is also why the
+            // device sort needs no cub::ThreadReduce, whose namespace shifted
+            // across CUB versions.)
         }
 
         __device__ __forceinline__