[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