From 8046b1a95df425b0b5c9e7f04c5ffe6f519d1802 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Fri, 6 Dec 2024 16:00:17 -0800 Subject: [PATCH 1/2] bug fix --- .../edge_partition_edge_property_device_view.cuh | 11 +++++++++-- .../edge_partition_endpoint_property_device_view.cuh | 11 +++++++++-- 2 files changed, 18 insertions(+), 4 deletions(-) diff --git a/cpp/include/cugraph/edge_partition_edge_property_device_view.cuh b/cpp/include/cugraph/edge_partition_edge_property_device_view.cuh index c7521831fb8..91a812351b2 100644 --- a/cpp/include/cugraph/edge_partition_edge_property_device_view.cuh +++ b/cpp/include/cugraph/edge_partition_edge_property_device_view.cuh @@ -138,9 +138,16 @@ class edge_partition_edge_property_device_view_t { { if constexpr (has_packed_bool_element) { static_assert(is_packed_bool, "unimplemented for thrust::tuple types."); + cuda::atomic_ref word( + *(value_first_ + cugraph::packed_bool_offset(offset))); auto mask = cugraph::packed_bool_mask(offset); - auto old = val ? atomicOr(value_first_ + cugraph::packed_bool_offset(offset), mask) - : atomicAnd(value_first_ + cugraph::packed_bool_offset(offset), ~mask); + uint32_t old{}; + if (compare == val) { + old = word.load(cuda::std::meory_order_relaxed); + } else { + old = val ? world.fetch_or(mask, cuda::std::meory_order_relaxed) + : world.fetch_and(~mask, cuda::std::meory_order_relaxed); + } return static_cast(old & mask); } else { return cugraph::elementwise_atomic_cas(value_first_ + offset, compare, val); diff --git a/cpp/include/cugraph/edge_partition_endpoint_property_device_view.cuh b/cpp/include/cugraph/edge_partition_endpoint_property_device_view.cuh index 824662a957e..7d65da3076e 100644 --- a/cpp/include/cugraph/edge_partition_endpoint_property_device_view.cuh +++ b/cpp/include/cugraph/edge_partition_endpoint_property_device_view.cuh @@ -145,9 +145,16 @@ class edge_partition_endpoint_property_device_view_t { auto val_offset = value_offset(offset); if constexpr (has_packed_bool_element) { static_assert(is_packed_bool, "unimplemented for thrust::tuple types."); + cuda::atomic_ref word( + *(value_first_ + cugraph::packed_bool_offset(val_offset))); auto mask = cugraph::packed_bool_mask(val_offset); - auto old = val ? atomicOr(value_first_ + cugraph::packed_bool_offset(val_offset), mask) - : atomicAnd(value_first_ + cugraph::packed_bool_offset(val_offset), ~mask); + uint32_t old{}; + if (compare == val) { + old = word.load(cuda::std::memory_order_relaxed); + } else { + old = val ? word.fetch_or(mask, cuda::std::memory_order_relaxed) + : word.fetch_and(~mask, cuda::std::memory_order_relaxed); + } return static_cast(old & mask); } else { return cugraph::elementwise_atomic_cas(value_first_ + val_offset, compare, val); From 18078e2331accf8db0e66161281106a54b9bb663 Mon Sep 17 00:00:00 2001 From: Seunghwa Kang Date: Fri, 6 Dec 2024 16:13:17 -0800 Subject: [PATCH 2/2] fix build error --- .../cugraph/edge_partition_edge_property_device_view.cuh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/include/cugraph/edge_partition_edge_property_device_view.cuh b/cpp/include/cugraph/edge_partition_edge_property_device_view.cuh index 91a812351b2..46cebbc122e 100644 --- a/cpp/include/cugraph/edge_partition_edge_property_device_view.cuh +++ b/cpp/include/cugraph/edge_partition_edge_property_device_view.cuh @@ -143,10 +143,10 @@ class edge_partition_edge_property_device_view_t { auto mask = cugraph::packed_bool_mask(offset); uint32_t old{}; if (compare == val) { - old = word.load(cuda::std::meory_order_relaxed); + old = word.load(cuda::std::memory_order_relaxed); } else { - old = val ? world.fetch_or(mask, cuda::std::meory_order_relaxed) - : world.fetch_and(~mask, cuda::std::meory_order_relaxed); + old = val ? word.fetch_or(mask, cuda::std::memory_order_relaxed) + : word.fetch_and(~mask, cuda::std::memory_order_relaxed); } return static_cast(old & mask); } else {