[Openmp-commits] [PATCH] D75587: [libomptarget][nfc] Move GetWarp/LaneId functions into per arch code
Jon Chesterfield via Phabricator via Openmp-commits
openmp-commits at lists.llvm.org
Tue Mar 3 17:16:48 PST 2020
JonChesterfield marked 2 inline comments as done.
JonChesterfield added inline comments.
================
Comment at: openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip:17
-static DEVICE uint32_t getLaneId(void) {
- return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
-}
-
// Initialized with a 64-bit mask with bits set in positions less than the
// thread's lane number in the warp
----------------
^ this is the preferred implementation, which is unfortunately different to the generic one
================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h:194
+INLINE unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); }
// Return true if this is the first active thread in the warp.
----------------
^nvptx uses exactly the same functions as before
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D75587/new/
https://reviews.llvm.org/D75587
More information about the Openmp-commits
mailing list