Fix "cudaErrorInvalidDevice: invalid device ordinal" in debug runs (#…

…4875)

This PR should replace PR #4872.

Louvain currently fails when built in the debug mode with an error message saying "cudaErrorInvalidDevice: invalid device ordinal".

This is due to a CUDA block size (set to 512) too large for debug runs due to resource limitations. This PR reduces the block sizes from 512 to 256.

Closes #4806 
Closes #4873

Authors:
  - Seunghwa Kang (https://github.com/seunghwak)
  - Chuck Hastings (https://github.com/ChuckHastings)

Approvers:
  - Joseph Nke (https://github.com/jnke2016)
  - Chuck Hastings (https://github.com/ChuckHastings)

URL: #4875

cpp/src/prims/detail/per_v_transform_reduce_e.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -69,7 +69,7 @@ namespace cugraph {
 
 namespace detail {
 
-int32_t constexpr per_v_transform_reduce_e_kernel_block_size                        = 512;
+int32_t constexpr per_v_transform_reduce_e_kernel_block_size                        = 256;
 int32_t constexpr per_v_transform_reduce_e_kernel_high_degree_reduce_any_block_size = 128;
 
 template <typename Iterator, typename default_t, typename Enable = void>
@@ -1610,7 +1610,7 @@ void per_v_transform_reduce_e(raft::handle_t const& handle,
         edge_partition.major_range_first(),
         handle.get_stream());
       assert((*key_segment_offsets).back() == *((*key_segment_offsets).rbegin() + 1));
-      assert(sorted_uniue_nzd_key_last == sorted_unique_key_first + (*key_segment_offsets).back());
+      assert(sorted_unique_nzd_key_last == sorted_unique_key_first + (*key_segment_offsets).back());
     }
   } else {
     tmp_vertex_value_output_first = vertex_value_output_first;

cpp/src/prims/fill_edge_src_dst_property.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -973,7 +973,7 @@ void fill_edge_minor_property(raft::handle_t const& handle,
     assert(graph_view.local_vertex_partition_range_size() ==
            (GraphViewType::is_storage_transposed
               ? graph_view.local_edge_partition_src_range_size()
-              : graph_view.local_edge_partition_dst_range_sizse()));
+              : graph_view.local_edge_partition_dst_range_size()));
     if constexpr (contains_packed_bool_element) {
       thrust::for_each(handle.get_thrust_policy(),
                        sorted_unique_vertex_first,

cpp/src/prims/vertex_frontier.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -227,8 +227,7 @@ void retrieve_vertex_list_from_bitmap(
 {
   using vertex_t = typename thrust::iterator_traits<OutputVertexIterator>::value_type;
 
-  assert((comm.get_rank() != root) ||
-         (bitmap.size() >= packed_bool_size(vertex_range_last - vertex_ragne_first)));
+  assert((bitmap.size() >= packed_bool_size(vertex_range_last - vertex_range_first)));
   detail::copy_if_nosync(thrust::make_counting_iterator(vertex_range_first),
                          thrust::make_counting_iterator(vertex_range_last),
                          thrust::make_transform_iterator(