From 481e22e34faaf55feab4263b9a39f12d53799d22 Mon Sep 17 00:00:00 2001 From: David Canar Date: Thu, 27 Aug 2026 13:16:59 -0600 Subject: [PATCH] hip: enable GGML_CUDA_USE_CUB via hipCUB Without CUB, ggml_cuda_op_top_k falls back to the bitonic argsort, which launches one block of next_power_of_2(ncols) threads. CUB was disabled unconditionally on HIP, so any ggml_top_k over more than 1024 rows (a GLM-5-Next / GLM-5.3-Flash lightning-indexer selection over more than 4096 KV cells, i.e. any conversation past ~4k tokens) asked for a >= 2048-thread block and aborted the backend with: ggml_cuda_compute_forward: TOP_K failed ROCm error: invalid configuration argument hipCUB provides every primitive the GGML_CUDA_USE_CUB paths use, so enable the define on HIP and bridge the differences in common.cuh: - alias namespace cub = hipcub (hipCUB has no ) - map cudaStreamCaptureStatus / cudaStreamIsCapturing onto the HIP names used by the capture checks in argsort.cu / mean.cu Two HIP-only adjustments inside the CUB paths: - keep the CCCL >= 3.2 DeviceTopK and CCCL >= 3.1 strided-iterator paths off on HIP (hipCUB has no cuda::execution / cuda/iterator) - argsort.cu: route the multi-row non-capturing path through DeviceSegmentedRadixSort instead of DeviceSegmentedSort, which hipCUB does not provide CUDA builds are untouched; everything new is inside GGML_USE_HIP. Verified on a Radeon 8060S (gfx1151) RPC pair running GLM-5.3-Flash UD-Q4_K_XL: an 8.8k-token prompt processes end to end where the bitonic path aborted the RPC server; prompt processing is unchanged. --- ggml/src/ggml-cuda/argsort.cu | 19 ++++++++++++++----- ggml/src/ggml-cuda/common.cuh | 17 +++++++++++++++-- ggml/src/ggml-cuda/cumsum.cu | 2 ++ ggml/src/ggml-cuda/mean.cu | 2 ++ ggml/src/ggml-cuda/sum.cu | 2 ++ ggml/src/ggml-cuda/top-k.cu | 4 +++- 6 files changed, 38 insertions(+), 8 deletions(-) diff --git a/ggml/src/ggml-cuda/argsort.cu b/ggml/src/ggml-cuda/argsort.cu index 26af9002597..b776e4bddc4 100644 --- a/ggml/src/ggml-cuda/argsort.cu +++ b/ggml/src/ggml-cuda/argsort.cu @@ -1,8 +1,10 @@ #include "argsort.cuh" #ifdef GGML_CUDA_USE_CUB +# ifndef GGML_USE_HIP // CUB comes from hipCUB via common.cuh on HIP # include -# if (CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 1) +# endif +# if !defined(GGML_USE_HIP) && (CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 1) # define STRIDED_ITERATOR_AVAILABLE # include # endif @@ -83,12 +85,19 @@ void argsort_f32_i32_cuda_cub(ggml_cuda_pool & pool, is_capturing = (capture_status != cudaStreamCaptureStatusNone); #endif // USE_CUDA_GRAPH +#ifdef GGML_USE_HIP + // hipCUB does not provide cub::DeviceSegmentedSort - use the radix segmented sort + const bool use_segmented_radix = true; +#else + const bool use_segmented_radix = is_capturing; +#endif + if (order == GGML_SORT_ORDER_ASC) { if (nrows == 1) { CUDA_CHECK(DeviceRadixSort::SortPairs(nullptr, temp_storage_bytes, temp_keys, temp_keys, // keys (in-place) temp_indices, dst, // values (indices) ncols, 0, sizeof(float) * 8, stream)); - } else if (is_capturing) { + } else if (use_segmented_radix) { CUDA_CHECK(DeviceSegmentedRadixSort::SortPairs( nullptr, temp_storage_bytes, temp_keys, temp_keys, // keys (in-place) temp_indices, dst, // values (indices) @@ -107,7 +116,7 @@ void argsort_f32_i32_cuda_cub(ggml_cuda_pool & pool, temp_keys, // keys (in-place) temp_indices, dst, // values (indices) ncols, 0, sizeof(float) * 8, stream)); - } else if (is_capturing) { + } else if (use_segmented_radix) { CUDA_CHECK(DeviceSegmentedRadixSort::SortPairsDescending( nullptr, temp_storage_bytes, temp_keys, temp_keys, temp_indices, dst, ncols * nrows, nrows, offset_iterator, offset_iterator + 1, 0, sizeof(float) * 8, stream)); @@ -127,7 +136,7 @@ void argsort_f32_i32_cuda_cub(ggml_cuda_pool & pool, temp_keys, // keys (in-place) temp_indices, dst, // values (indices) ncols, 0, sizeof(float) * 8, stream)); - } else if (is_capturing) { + } else if (use_segmented_radix) { CUDA_CHECK(DeviceSegmentedRadixSort::SortPairs(d_temp_storage, temp_storage_bytes, temp_keys, temp_keys, temp_indices, dst, ncols * nrows, nrows, offset_iterator, offset_iterator + 1, 0, sizeof(float) * 8, stream)); @@ -142,7 +151,7 @@ void argsort_f32_i32_cuda_cub(ggml_cuda_pool & pool, temp_keys, // keys (in-place) temp_indices, dst, // values (indices) ncols, 0, sizeof(float) * 8, stream)); - } else if (is_capturing) { + } else if (use_segmented_radix) { CUDA_CHECK(DeviceSegmentedRadixSort::SortPairsDescending( d_temp_storage, temp_storage_bytes, temp_keys, temp_keys, temp_indices, dst, ncols * nrows, nrows, offset_iterator, offset_iterator + 1, 0, sizeof(float) * 8, stream)); diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 14dd1098c97..0d432470f46 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -107,9 +107,22 @@ #define GGML_CUDA_CC_IS_QY2(cc) (cc >= GGML_CUDA_CC_QY2 && cc < GGML_CUDA_CC_PH1) #define GGML_CUDA_CC_IS_PH1(cc) (cc >= GGML_CUDA_CC_PH1) -#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070 +#if (!defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070) || defined(GGML_USE_HIP) # define GGML_CUDA_USE_CUB -#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070 +#endif // (!defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070) || defined(GGML_USE_HIP) + +#ifdef GGML_USE_HIP +// On HIP builds CUB comes from hipCUB: the primitives live in namespace hipcub and +// there is no to include, so alias the namespace and map the few CUDA +// stream-capture names the GGML_CUDA_USE_CUB code paths use onto the HIP ones. +# include +namespace cub = hipcub; +using cudaStreamCaptureStatus = hipStreamCaptureStatus; +constexpr auto cudaStreamCaptureStatusNone = hipStreamCaptureStatusNone; +inline cudaError_t cudaStreamIsCapturing(cudaStream_t stream, cudaStreamCaptureStatus * status) { + return hipStreamIsCapturing(stream, status); +} +#endif // GGML_USE_HIP // PDL host-side support (cudaLaunchKernelEx) requires CUDART >= 11.8. // However, this has been bugged in CTK < 12.3 for MSVC builds, see diff --git a/ggml/src/ggml-cuda/cumsum.cu b/ggml/src/ggml-cuda/cumsum.cu index def9c32955f..327f5eeb470 100644 --- a/ggml/src/ggml-cuda/cumsum.cu +++ b/ggml/src/ggml-cuda/cumsum.cu @@ -5,7 +5,9 @@ #include "ggml.h" #ifdef GGML_CUDA_USE_CUB +# ifndef GGML_USE_HIP // CUB comes from hipCUB via common.cuh on HIP # include +# endif #endif // GGML_CUDA_USE_CUB template diff --git a/ggml/src/ggml-cuda/mean.cu b/ggml/src/ggml-cuda/mean.cu index a8f6046e46d..87afc1ab3dd 100644 --- a/ggml/src/ggml-cuda/mean.cu +++ b/ggml/src/ggml-cuda/mean.cu @@ -2,7 +2,9 @@ #include "reduce_rows.cuh" #ifdef GGML_CUDA_USE_CUB +# ifndef GGML_USE_HIP // CUB comes from hipCUB via common.cuh on HIP #include +# endif using namespace cub; #endif // GGML_CUDA_USE_CUB diff --git a/ggml/src/ggml-cuda/sum.cu b/ggml/src/ggml-cuda/sum.cu index c56257b4406..69a9ac554c0 100644 --- a/ggml/src/ggml-cuda/sum.cu +++ b/ggml/src/ggml-cuda/sum.cu @@ -2,7 +2,9 @@ #include "sumrows.cuh" #ifdef GGML_CUDA_USE_CUB +# ifndef GGML_USE_HIP // CUB comes from hipCUB via common.cuh on HIP #include +# endif using namespace cub; #endif // GGML_CUDA_USE_CUB diff --git a/ggml/src/ggml-cuda/top-k.cu b/ggml/src/ggml-cuda/top-k.cu index 9681cd29333..b7ca37e00b2 100644 --- a/ggml/src/ggml-cuda/top-k.cu +++ b/ggml/src/ggml-cuda/top-k.cu @@ -2,8 +2,10 @@ #include "top-k.cuh" #ifdef GGML_CUDA_USE_CUB +# ifndef GGML_USE_HIP // CUB comes from hipCUB via common.cuh on HIP # include -# if (CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 2) +# endif +# if !defined(GGML_USE_HIP) && (CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 2) # define CUB_TOP_K_AVAILABLE # include using namespace cub;