[Openmp-commits] [PATCH] D75587: [libomptarget][nfc] Move GetWarp/LaneId functions into per arch code

[Jon Chesterfield via Phabricator via Openmp-commits]

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.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D75587

Files:
  openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
  openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip
  openmp/libomptarget/deviceRTLs/common/src/support.cu
  openmp/libomptarget/deviceRTLs/common/support.h
  openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h

-------------- 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>