diff options
| author | George Rokos <grokos@us.ibm.com> | 2018-03-14 19:11:36 +0000 |
|---|---|---|
| committer | George Rokos <grokos@us.ibm.com> | 2018-03-14 19:11:36 +0000 |
| commit | 59be4b434f5629443bf26e86d261ba618851b31b (patch) | |
| tree | b7aebc4f88f3b8524402d6b3f2609b6d78c38b46 | |
| parent | 4e6b822cdc2b4780c084aff39f59bd60cba0260c (diff) | |
| download | bcm5719-llvm-59be4b434f5629443bf26e86d261ba618851b31b.tar.gz bcm5719-llvm-59be4b434f5629443bf26e86d261ba618851b31b.zip | |
[libomptarget][nvptx] Bug fix: Correctly identify the warp master active thread.
llvm-svn: 327556
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu | 3 |
1 files changed, 2 insertions, 1 deletions
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu index e2a38e35e0b..cd73a6b4108 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() { |
