From: George Rokos <grokos@us.ibm.com>
Date: Wed, 14 Mar 2018 19:11:36 +0000 (+0000)
Subject: [libomptarget][nvptx] Bug fix: Correctly identify the warp master active thread.
X-Git-Tag: llvmorg-7.0.0-rc1~10576
X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=59be4b434f5629443bf26e86d261ba618851b31b;p=platform%2Fupstream%2Fllvm.git

[libomptarget][nvptx] Bug fix: Correctly identify the warp master active thread.

llvm-svn: 327556
---

diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
index e2a38e3..cd73a6b 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
@@ -37,7 +37,8 @@ __device__ static bool IsWarpMasterActiveThread() {
   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() {