ggml-cuda: missing <cuda/iterator>/<cuda/execution>/<cuda/stream_ref> includes break build with CCCL 3.1+ (argsort.cu, top-k.cu)

Description

Building ggml-cuda fails with CUDA Toolkit 13.4 (CCCL 3.4.2 bundled) because two files use cuda:: (libcu++/CCCL) iterator and execution-policy APIs without including the headers that declare them. Only <cub/cub.cuh> is included, which does not pull in <cuda/iterator>, <cuda/execution>, or <cuda/stream_ref>.

This is not a toolkit compatibility issue — the required CCCL headers are present in the CUDA 13.4 install and provide the needed APIs; they are simply never #included.

Affected files

ggml/src/ggml-cuda/argsort.cu

#ifdef GGML_CUDA_USE_CUB
#    include <cub/cub.cuh>
#    if (CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 1)
#        define STRIDED_ITERATOR_AVAILABLE
#    endif
using namespace cub;
#endif  // GGML_CUDA_USE_CUB

STRIDED_ITERATOR_AVAILABLE gates use of cuda::make_strided_iterator(cuda::make_counting_iterator(0), ncols) later in the file, but neither cuda::make_strided_iterator nor cuda::make_counting_iterator are declared by <cub/cub.cuh>.

ggml/src/ggml-cuda/top-k.cu

#ifdef GGML_CUDA_USE_CUB
#    include <cub/cub.cuh>
#    if (CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 2)
#        define CUB_TOP_K_AVAILABLE
using namespace cub;
#    endif  // CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 2
#endif      // GGML_CUDA_USE_CUB

CUB_TOP_K_AVAILABLE gates code using cuda::execution::require, cuda::execution::determinism, cuda::execution::output_ordering, cuda::stream_ref, cuda::std::execution::env, cuda::make_counting_iterator, and cuda::discard_iterator — none of which are declared by <cub/cub.cuh> either.

Error log

argsort.cu(48): error: namespace "cuda" has no member "make_strided_iterator"
      auto offset_iterator = cuda::make_strided_iterator(cuda::make_counting_iterator(0), ncols);
                                   ^

argsort.cu(48): error: namespace "cuda" has no member "make_counting_iterator"
      auto offset_iterator = cuda::make_strided_iterator(cuda::make_counting_iterator(0), ncols);
                                                               ^

top-k.cu(25): error: namespace "cuda" has no member "make_counting_iterator"
top-k.cu(28): error: incomplete type "cuda::__4::discard_iterator" is not allowed
top-k.cu(28): error: no instance of overloaded function "cub::_V_300402_SM_1200::DeviceTopK::MaxPairs" matches the argument list

Fix

Add the missing includes inside the version-gated blocks:

argsort.cu:

#    if (CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 1)
#        define STRIDED_ITERATOR_AVAILABLE
#        include <cuda/iterator>
#    endif

top-k.cu:

#    if (CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 2)
#        define CUB_TOP_K_AVAILABLE
#        include <cuda/execution>
#        include <cuda/iterator>
#        include <cuda/stream_ref>
using namespace cub;
#    endif

With these includes added, both files compile cleanly (one deprecation warning from <cuda/stream_ref>, which upstream CCCL says will move to <cuda/stream> in a future release — worth switching to <cuda/stream> directly if it's already available in supported CCCL versions).

Environment

• OS: Windows 11
• Compiler: MSVC (Visual Studio 2022, via vcvars64.bat), nvcc from CUDA 13.4
• CUDA Toolkit: 13.4.59
• CCCL version bundled with toolkit: 3.4.2 (CCCL_VERSION 3004002)
• GPU: NVIDIA GeForce RTX 5060 Ti (Blackwell, sm_120a)
• Build system: CMake + Ninja
• ggml build type used: GGML_CUDA=ON, CUB path enabled (GGML_CUDA_USE_CUB)

Steps to reproduce

1. Build ggml with -DGGML_CUDA=ON targeting an architecture that enables GGML_CUDA_USE_CUB (any recent GPU) using CUDA Toolkit 13.4 (CCCL 3.4.2).
2. ninja fails compiling argsort.cu.obj and top-k.cu.obj with the errors above.