From 638f84deed8ea662dbed3853ec21c76adc4bbf6b Mon Sep 17 00:00:00 2001
From: Chris Maffeo <cmaffeo2@illinois.edu>
Date: Fri, 31 Jul 2020 13:04:43 -0500
Subject: [PATCH] Reverted from _ballot_sync() to __activemask() in
atomicAggInc
---
src/CudaUtil.cu | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/src/CudaUtil.cu b/src/CudaUtil.cu
index eb23592..5068fe8 100644
--- a/src/CudaUtil.cu
+++ b/src/CudaUtil.cu
@@ -41,7 +41,8 @@ __inline__ __device__ uint __lanemask_lt()
}
__device__ int atomicAggInc(int *ctr, int warpLane)
{
- unsigned int active = __ballot_sync(0xFFFFFFFF, 1);
+ // unsigned int active = __ballot_sync(0xFFFFFFFF, 1);
+ unsigned int active = __activemask();
int leader = __ffs(active) - 1;
int change = __popc(active);
unsigned int rank = __popc(active & __lanemask_lt());
--
GitLab