[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 created this revision.
JonChesterfield added reviewers: jdoerfert, grokos, ABataev.
Herald added subscribers: openmp-commits, jvesely.
Herald added a project: OpenMP.
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

[libomptarget][nfc] Move GetWarp/LaneId functions into per arch code

No code change for nvptx. Amdgcn currently has two implementations of GetLaneId,
this patch keeps the one a colleague considered to be superior for our ISA.

GetWarpId is currently the same function for amdgcn and nvptx, but I think it's
cleaner to keep it grouped with all the others than to keep it in support.cu.

  rG LLVM Github Monorepo



-------------- next part --------------
A non-text attachment was scrubbed...
Name: D75587.248074.patch
Type: text/x-patch
Size: 5106 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20200304/3333a0bf/attachment-0001.bin>

More information about the Openmp-commits mailing list