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

Add patch for thrust-cub 1.16 to fix sort compile times #10577

Merged
merged 1 commit into from
Apr 4, 2022
Merged

Add patch for thrust-cub 1.16 to fix sort compile times #10577

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
102 changes: 53 additions & 49 deletions cpp/cmake/thrust.patch
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,52 +1,39 @@
diff --git a/thrust/system/cuda/detail/sort.h b/thrust/system/cuda/detail/sort.h
index 1ffeef0..5e80800 100644
--- a/thrust/system/cuda/detail/sort.h
+++ b/thrust/system/cuda/detail/sort.h
@@ -108,7 +108,7 @@ namespace __merge_sort {
key_type key2 = keys_shared[keys2_beg];

diff --git a/cub/block/block_merge_sort.cuh b/cub/block/block_merge_sort.cuh
index 4769df36..d86d6342 100644
--- a/cub/block/block_merge_sort.cuh
+++ b/cub/block/block_merge_sort.cuh
@@ -91,7 +91,7 @@ __device__ __forceinline__ void SerialMerge(KeyT *keys_shared,
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) &&
@@ -311,10 +311,10 @@ namespace __merge_sort {
void stable_odd_even_sort(key_type (&keys)[ITEMS_PER_THREAD],
item_type (&items)[ITEMS_PER_THREAD])
for (int item = 0; item < ITEMS_PER_THREAD; ++item)
{
bool p = (keys2_beg < keys2_end) &&
@@ -383,7 +383,7 @@ public:
//
KeyT max_key = oob_default;

- #pragma unroll
+ #pragma unroll 1
for (int item = 1; item < ITEMS_PER_THREAD; ++item)
{
-#pragma unroll
+#pragma unroll 1
for (int i = 0; i < ITEMS_PER_THREAD; ++i)
{
-#pragma unroll
+#pragma unroll 1
for (int j = 1 & i; j < ITEMS_PER_THREAD - 1; j += 2)
{
if (compare_op(keys[j + 1], keys[j]))
@@ -350,7 +350,7 @@ namespace __merge_sort {
// each thread has sorted keys_loc
// merge sort keys_loc in shared memory
//
-#pragma unroll
+#pragma unroll 1
for (int coop = 2; coop <= BLOCK_THREADS; coop *= 2)
{
sync_threadblock();
@@ -479,7 +479,7 @@ namespace __merge_sort {
// and fill the remainig keys with it
//
key_type max_key = keys_loc[0];
-#pragma unroll
+#pragma unroll 1
for (int ITEM = 1; ITEM < ITEMS_PER_THREAD; ++ITEM)
{
if (ITEMS_PER_THREAD * tid + ITEM < num_remaining)
diff a/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/device/dispatch/dispatch_radix_sort.cuh
index 41eb1d2..f2893b4 100644
if (ITEMS_PER_THREAD * linear_tid + item < valid_items)
@@ -407,7 +407,7 @@ public:
// each thread has sorted keys
// merge sort keys in shared memory
//
- #pragma unroll
+ #pragma unroll 1
for (int target_merged_threads_number = 2;
target_merged_threads_number <= NUM_THREADS;
target_merged_threads_number *= 2)
diff --git a/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/device/dispatch/dispatch_radix_sort.cuh
index b188c75f..3f36656f 100644
--- a/cub/device/dispatch/dispatch_radix_sort.cuh
+++ b/cub/device/dispatch/dispatch_radix_sort.cuh
@@ -723,7 +723,7 @@ struct DeviceRadixSortPolicy
@@ -736,7 +736,7 @@ struct DeviceRadixSortPolicy


/// SM60 (GP100)
Expand All @@ -55,11 +42,11 @@ index 41eb1d2..f2893b4 100644
{
enum {
PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5, // 6.9B 32b keys/s (Quadro P100)
diff a/cub/device/dispatch/dispatch_reduce.cuh b/cub/device/dispatch/dispatch_reduce.cuh
index f6aee45..dd64301 100644
diff --git a/cub/device/dispatch/dispatch_reduce.cuh b/cub/device/dispatch/dispatch_reduce.cuh
index e0470ccb..6a0c2ed6 100644
--- a/cub/device/dispatch/dispatch_reduce.cuh
+++ b/cub/device/dispatch/dispatch_reduce.cuh
@@ -284,7 +284,7 @@ struct DeviceReducePolicy
@@ -280,7 +280,7 @@ struct DeviceReducePolicy
};

/// SM60
Expand All @@ -68,11 +55,11 @@ index f6aee45..dd64301 100644
{
// ReducePolicy (P100: 591 GB/s @ 64M 4B items; 583 GB/s @ 256M 1B items)
typedef AgentReducePolicy<
diff a/cub/device/dispatch/dispatch_scan.cuh b/cub/device/dispatch/dispatch_scan.cuh
index c0c6d59..937ee31 100644
diff --git a/cub/device/dispatch/dispatch_scan.cuh b/cub/device/dispatch/dispatch_scan.cuh
index c2d04588..ac2d10e0 100644
--- a/cub/device/dispatch/dispatch_scan.cuh
+++ b/cub/device/dispatch/dispatch_scan.cuh
@@ -178,7 +178,7 @@ struct DeviceScanPolicy
@@ -177,7 +177,7 @@ struct DeviceScanPolicy
};

/// SM600
Expand All @@ -81,3 +68,20 @@ index c0c6d59..937ee31 100644
{
typedef AgentScanPolicy<
128, 15, ///< Threads per block, items per thread
diff --git a/cub/thread/thread_sort.cuh b/cub/thread/thread_sort.cuh
index 5d486789..b42fb5f0 100644
--- a/cub/thread/thread_sort.cuh
+++ b/cub/thread/thread_sort.cuh
@@ -83,10 +83,10 @@ StableOddEvenSort(KeyT (&keys)[ITEMS_PER_THREAD],
{
constexpr bool KEYS_ONLY = std::is_same<ValueT, NullType>::value;

- #pragma unroll
+ #pragma unroll 1
for (int i = 0; i < ITEMS_PER_THREAD; ++i)
{
- #pragma unroll
+ #pragma unroll 1
for (int j = 1 & i; j < ITEMS_PER_THREAD - 1; j += 2)
{
if (compare_op(keys[j + 1], keys[j]))