Merge branch 'branch-25.04' of https://github.com/rapidsai/cugraph in…

…to fea_find_forest2
rapidsai · Mar 6, 2025 · ff00225 · ff00225
2 parents 09f54fc + 6618212
commit ff00225
Showing 1 changed file with 18 additions and 4 deletions.
diff --git a/cpp/src/prims/detail/per_v_transform_reduce_e.cuh b/cpp/src/prims/detail/per_v_transform_reduce_e.cuh
@@ -589,6 +589,13 @@ __global__ static void per_v_transform_reduce_e_mid_degree(
       }
     }
 
+    // FIXME: Remove once upgraded to CCCL version 3.x
+#if CCCL_MAJOR_VERSION >= 3
+    using cuda::minimum;
+#else
+    using minimum = cub::Min;
+#endif
+
     if (edge_partition_e_mask) {
       if constexpr (update_major && std::is_same_v<ReduceOp, reduce_op::any<T>>) {
         auto rounded_up_local_degree =
@@ -601,7 +608,7 @@ __global__ static void per_v_transform_reduce_e_mid_degree(
             e_op_result = call_e_op(i);
           }
           first_valid_lane_id = WarpReduce(temp_storage[threadIdx.x / raft::warp_size()])
-                                  .Reduce(e_op_result ? lane_id : raft::warp_size(), cub::Min());
+                                  .Reduce(e_op_result ? lane_id : raft::warp_size(), minimum{});
           first_valid_lane_id = __shfl_sync(raft::warp_full_mask(), first_valid_lane_id, int{0});
           if (lane_id == first_valid_lane_id) { reduced_e_op_result = *e_op_result; }
           if (first_valid_lane_id != raft::warp_size()) { break; }
@@ -634,7 +641,7 @@ __global__ static void per_v_transform_reduce_e_mid_degree(
             e_op_result = call_e_op(i);
           }
           first_valid_lane_id = WarpReduce(temp_storage[threadIdx.x / raft::warp_size()])
-                                  .Reduce(e_op_result ? lane_id : raft::warp_size(), cub::Min());
+                                  .Reduce(e_op_result ? lane_id : raft::warp_size(), minimum{});
           first_valid_lane_id = __shfl_sync(raft::warp_full_mask(), first_valid_lane_id, int{0});
           if (lane_id == first_valid_lane_id) { reduced_e_op_result = *e_op_result; }
           if (first_valid_lane_id != raft::warp_size()) { break; }
@@ -781,6 +788,13 @@ __global__ static void per_v_transform_reduce_e_high_degree(
       }
     }
 
+    // FIXME: Remove once upgraded to CCCL version 3.x
+#if CCCL_MAJOR_VERSION >= 3
+    using cuda::minimum;
+#else
+    using minimum = cub::Min;
+#endif
+
     if (edge_partition_e_mask) {
       if constexpr (update_major && std::is_same_v<ReduceOp, reduce_op::any<T>>) {
         auto rounded_up_local_degree =
@@ -799,7 +813,7 @@ __global__ static void per_v_transform_reduce_e_high_degree(
               .Reduce(e_op_result
                         ? threadIdx.x
                         : per_v_transform_reduce_e_kernel_high_degree_reduce_any_block_size,
-                      cub::Min());
+                      minimum{});
           if (threadIdx.x == 0) { output_thread_id = first_valid_thread_id; }
           __syncthreads();
           first_valid_thread_id = output_thread_id;
@@ -843,7 +857,7 @@ __global__ static void per_v_transform_reduce_e_high_degree(
               .Reduce(e_op_result
                         ? threadIdx.x
                         : per_v_transform_reduce_e_kernel_high_degree_reduce_any_block_size,
-                      cub::Min());
+                      minimum{});
           if (threadIdx.x == 0) { output_thread_id = first_valid_thread_id; }
           __syncthreads();
           first_valid_thread_id = output_thread_id;