Skip to content
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

Jump to bottom

Update patches for CCCL 2.6 #16668

Open
wants to merge 1 commit into
base: branch-24.10
Choose a base branch
from
+56 −22
Open

Update patches for CCCL 2.6 #16668

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
66 changes: 50 additions & 16 deletions cpp/cmake/thirdparty/patches/thrust_disable_64bit_dispatching.diff
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,25 +1,59 @@
diff --git a/thrust/thrust/system/cuda/detail/dispatch.h b/thrust/thrust/system/cuda/detail/dispatch.h
index 2a3cc4e33..8fb337b26 100644
index 971b93d62..0d6b25b07 100644
--- a/thrust/thrust/system/cuda/detail/dispatch.h
+++ b/thrust/thrust/system/cuda/detail/dispatch.h
@@ -44,8 +44,7 @@
} \
else \
{ \
- auto THRUST_PP_CAT2(count, _fixed) = static_cast<thrust::detail::int64_t>(count); \
- status = call arguments; \
@@ -36,16 +36,15 @@
* that callables for both branches consist of the same tokens, and is intended to be used with Thrust-style dispatch
* interfaces, that always deduce the size type from the arguments.
*/
-#define THRUST_INDEX_TYPE_DISPATCH(status, call, count, arguments) \
- if (count <= thrust::detail::integer_traits<std::int32_t>::const_max) \
- { \
- auto THRUST_PP_CAT2(count, _fixed) = static_cast<std::int32_t>(count); \
- status = call arguments; \
- } \
- else \
- { \
- auto THRUST_PP_CAT2(count, _fixed) = static_cast<std::int64_t>(count); \
- status = call arguments; \
+#define THRUST_INDEX_TYPE_DISPATCH(status, call, count, arguments) \
+ if (count <= thrust::detail::integer_traits<std::int32_t>::const_max) \
+ { \
+ auto THRUST_PP_CAT2(count, _fixed) = static_cast<std::int32_t>(count); \
+ status = call arguments; \
+ } \
+ else \
+ { \
+ throw std::runtime_error("THRUST_INDEX_TYPE_DISPATCH 64-bit count is unsupported in libcudf"); \
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Our formatter did some crazy things here, so that the patches are less localized. Shout if you need them to be minimal

}

/**
@@ -66,9 +65,7 @@
} \
else \
{ \
- auto THRUST_PP_CAT2(count1, _fixed) = static_cast<thrust::detail::int64_t>(count1); \
- auto THRUST_PP_CAT2(count2, _fixed) = static_cast<thrust::detail::int64_t>(count2); \
- status = call arguments; \
@@ -55,18 +54,16 @@
*
* This version of the macro supports providing two count variables, which is necessary for set algorithms.
*/
-#define THRUST_DOUBLE_INDEX_TYPE_DISPATCH(status, call, count1, count2, arguments) \
- if (count1 + count2 <= thrust::detail::integer_traits<std::int32_t>::const_max) \
- { \
- auto THRUST_PP_CAT2(count1, _fixed) = static_cast<std::int32_t>(count1); \
- auto THRUST_PP_CAT2(count2, _fixed) = static_cast<std::int32_t>(count2); \
- status = call arguments; \
- } \
- else \
- { \
- auto THRUST_PP_CAT2(count1, _fixed) = static_cast<std::int64_t>(count1); \
- auto THRUST_PP_CAT2(count2, _fixed) = static_cast<std::int64_t>(count2); \
- status = call arguments; \
+#define THRUST_DOUBLE_INDEX_TYPE_DISPATCH(status, call, count1, count2, arguments) \
+ if (count1 + count2 <= thrust::detail::integer_traits<std::int32_t>::const_max) \
+ { \
+ auto THRUST_PP_CAT2(count1, _fixed) = static_cast<std::int32_t>(count1); \
+ auto THRUST_PP_CAT2(count2, _fixed) = static_cast<std::int32_t>(count2); \
+ status = call arguments; \
+ } \
+ else \
+ { \
+ throw std::runtime_error("THRUST_DOUBLE_INDEX_TYPE_DISPATCH 64-bit count is unsupported in libcudf"); \
}

/**
* Dispatch between 32-bit and 64-bit index based versions of the same algorithm
12 changes: 6 additions & 6 deletions cpp/cmake/thirdparty/patches/thrust_faster_sort_compile_times.diff
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,20 +1,20 @@
diff --git a/cub/cub/block/block_merge_sort.cuh b/cub/cub/block/block_merge_sort.cuh
index eb76ebb0b..c6c529a50 100644
index 29510db5e..cf57e5786 100644
--- a/cub/cub/block/block_merge_sort.cuh
+++ b/cub/cub/block/block_merge_sort.cuh
@@ -95,7 +95,7 @@ _CCCL_DEVICE _CCCL_FORCEINLINE void SerialMerge(
KeyT key1 = keys_shared[keys1_beg];
KeyT key2 = keys_shared[keys2_beg];

-#pragma unroll
+#pragma unroll 1
for (int item = 0; item < ITEMS_PER_THREAD; ++item)
{
bool p = (keys2_beg < keys2_end) && ((keys1_beg >= keys1_end) || compare_op(key2, key1));
@@ -376,7 +376,7 @@ public:
const bool p = (keys2_beg < keys2_end) && ((keys1_beg >= keys1_end) || compare_op(key2, key1));
@@ -374,7 +374,7 @@ public:
//
KeyT max_key = oob_default;

-#pragma unroll
+#pragma unroll 1
for (int item = 1; item < ITEMS_PER_THREAD; ++item)
Expand All @@ -27,7 +27,7 @@ index 7d9e8622f..da5627306 100644
@@ -87,10 +87,10 @@ StableOddEvenSort(KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&items)[ITEMS_PER_THRE
{
constexpr bool KEYS_ONLY = ::cuda::std::is_same<ValueT, NullType>::value;

-#pragma unroll
+#pragma unroll 1
for (int i = 0; i < ITEMS_PER_THREAD; ++i)
Expand Down
Loading