Skip to content

hip: enable GGML_CUDA_USE_CUB via hipCUB (fixes TOP_K abort for >1024 rows) #136

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
davidcanar wants to merge 1 commit into
base: glm5next/upstream
Choose a base branch
from
Open

hip: enable GGML_CUDA_USE_CUB via hipCUB (fixes TOP_K abort for >1024 rows) #136

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 14 additions & 5 deletions ggml/src/ggml-cuda/argsort.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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 <cub/cub.cuh>
# 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 <cuda/iterator>
# endif
Expand Down Expand Up @@ -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)
Expand All @@ -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));
Expand All @@ -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));
Expand All @@ -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));
Expand Down
17 changes: 15 additions & 2 deletions ggml/src/ggml-cuda/common.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <cub/cub.cuh> 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 <hipcub/hipcub.hpp>
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
Expand Down
2 changes: 2 additions & 0 deletions ggml/src/ggml-cuda/cumsum.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <cub/cub.cuh>
# endif
#endif // GGML_CUDA_USE_CUB

template<typename T, int BLOCK_SIZE>
Expand Down
2 changes: 2 additions & 0 deletions ggml/src/ggml-cuda/mean.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <cub/cub.cuh>
# endif
using namespace cub;
#endif // GGML_CUDA_USE_CUB

Expand Down
2 changes: 2 additions & 0 deletions ggml/src/ggml-cuda/sum.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <cub/cub.cuh>
# endif
using namespace cub;
#endif // GGML_CUDA_USE_CUB

Expand Down
4 changes: 3 additions & 1 deletion ggml/src/ggml-cuda/top-k.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <cub/cub.cuh>
# 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 <cuda/iterator>
using namespace cub;
Expand Down