diff --git a/pyproject.toml b/pyproject.toml index c5925b4b..12294ffd 100644 --- a/pyproject.toml +++ b/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 diff --git a/src/bridge/CMakeLists.txt b/src/bridge/CMakeLists.txt index b3b44de6..07ade7ec 100644 --- a/src/bridge/CMakeLists.txt +++ b/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() diff --git a/src/mw/device/sort_archetype.cpp b/src/mw/device/sort_archetype.cpp index 3f758a32..a8c3b73f 100644 --- a/src/mw/device/sort_archetype.cpp +++ b/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__