[Openmp-commits] [PATCH] D108398: [libomptarget] Specialize amdgpu devicertl on wave size for gfx10

Jon Chesterfield via Phabricator via Openmp-commits openmp-commits at lists.llvm.org
Thu Aug 26 02:12:53 PDT 2021


JonChesterfield planned changes to this revision.
JonChesterfield added a comment.

Only useful after D108708 <https://reviews.llvm.org/D108708> has landed



================
Comment at: openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip:145
 EXTERN unsigned GetLaneId() {
-  return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
+#if __AMDGCN_WAVEFRONT_SIZE == 64
+    return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
----------------
This should be a 'runtime' switch on gridvalues, will update the patch


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D108398/new/

https://reviews.llvm.org/D108398



More information about the Openmp-commits mailing list