[Openmp-commits] [PATCH] D88602: [libomptarget][amdgcn] Implement partial barrier

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield created this revision.
JonChesterfield added reviewers: jdoerfert, ABataev, grokos, ye-luo, tianshilei1992, jhuber6.
Herald added subscribers: openmp-commits, jfb, jvesely.
Herald added a project: OpenMP.
JonChesterfield requested review of this revision.
Herald added a subscriber: sstefan1.

[libomptarget][amdgcn] Implement partial barrier

named_sync is used to coordinate non-spmd kernels. This uses bar.sync on nvptx.
There is no corresponding ISA support on amdgcn, so this is implemented using
shared memory, one word initialized to zero.

Each wave increments the variable by one. Whichever wave is last is responsible
for resetting the variable to zero, at which point it and the others continue.

The race condition on a wave reaching the barrier before another wave has
noticed that it has been released is handled with a generation counter, packed
into the same word.

Uses a shared variable that is not needed on nvptx. Introduces a new hook,
kmpc_impl_target_init, to allow different targets to do extra initialization.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D88602

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

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D88602.295367.patch
Type: text/x-patch
Size: 4348 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20200930/2f8d437a/attachment.bin>