reduce overhead of warp-sized cooperative groups

ginkgo-project · Jan 13, 2020 · 1771533 · 1771533
1 parent b1fec39
commit 1771533
Showing 1 changed file with 20 additions and 4 deletions.
diff --git a/hip/components/cooperative_groups.hip.hpp b/hip/components/cooperative_groups.hip.hpp
@@ -250,7 +250,11 @@ class thread_block_tile {
     __device__ __forceinline__ int any(int predicate) const noexcept
     {
 #if GINKGO_HIP_PLATFORM_HCC
-        return (__ballot(predicate) & data_.mask) != 0;
+        if (Size == config::warp_size) {
+            return __any(predicate);
+        } else {
+            return (__ballot(predicate) & data_.mask) != 0;
+        }
 #else
         return __any_sync(data_.mask, predicate);
 #endif
@@ -259,7 +263,11 @@ class thread_block_tile {
     __device__ __forceinline__ int all(int predicate) const noexcept
     {
 #if GINKGO_HIP_PLATFORM_HCC
-        return (__ballot(predicate) & data_.mask) == data_.mask;
+        if (Size == config::warp_size) {
+            return __all(predicate);
+        } else {
+            return (__ballot(predicate) & data_.mask) == data_.mask;
+        }
 #else
         return __all_sync(data_.mask, predicate);
 #endif
@@ -269,9 +277,17 @@ class thread_block_tile {
         int predicate) const noexcept
     {
 #if GINKGO_HIP_PLATFORM_HCC
-        return (__ballot(predicate) & data_.mask) >> data_.lane_offset;
+        if (Size == config::warp_size) {
+            return __ballot(predicate);
+        } else {
+            return (__ballot(predicate) & data_.mask) >> data_.lane_offset;
+        }
 #else
-        return __ballot_sync(data_.mask, predicate) >> data_.lane_offset;
+        if (Size == config::warp_size) {
+            return __ballot_sync(data_.mask, predicate);
+        } else {
+            return __ballot_sync(data_.mask, predicate) >> data_.lane_offset;
+        }
 #endif
     }