[Openmp-commits] [openmp] r327556 - [libomptarget][nvptx] Bug fix: Correctly identify the warp master active thread.

[George Rokos via Openmp-commits]

Author: grokos
Date: Wed Mar 14 12:11:36 2018
New Revision: 327556

URL: http://llvm.org/viewvc/llvm-project?rev=327556&view=rev
Log:
[libomptarget][nvptx] Bug fix: Correctly identify the warp master active thread.


Modified:
    openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu?rev=327556&r1=327555&r2=327556&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu Wed Mar 14 12:11:36 2018
@@ -37,7 +37,8 @@ __device__ static bool IsWarpMasterActiv
   unsigned long long Mask = getActiveThreadsMask();
   unsigned long long ShNum = WARPSIZE - (getThreadId() % WARPSIZE);
   unsigned long long Sh = Mask << ShNum;
-  return Sh == 0;
+  // Truncate Sh to the 32 lower bits
+  return (unsigned)Sh == 0;
 }
 // Return true if this is the master thread.
 __device__ static bool IsMasterThread() {