[all-commits] [llvm/llvm-project] 8b6cd1: [libomptarget][amdgcn] Implement partial barrier

[Jon Chesterfield via All-commits]

  Branch: refs/heads/master
  Home:   https://github.com/llvm/llvm-project
  Commit: 8b6cd15242673c04618fb0aafc07d5de9e0bbe1e
      https://github.com/llvm/llvm-project/commit/8b6cd15242673c04618fb0aafc07d5de9e0bbe1e
  Author: JonChesterfield <jonathanchesterfield at gmail.com>
  Date:   2020-10-12 (Mon, 12 Oct 2020)

  Changed paths:
    M openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
    M openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip
    M openmp/libomptarget/deviceRTLs/common/src/omptarget.cu
    M openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h

  Log Message:
  -----------
  [libomptarget][amdgcn] Implement partial barrier

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

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D88602