[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