[all-commits] [llvm/llvm-project] d25679: [nfc][libomptarget] Drop parameter to named_sync

Jon Chesterfield via All-commits all-commits at lists.llvm.org
Tue Sep 29 15:12:51 PDT 2020


  Branch: refs/heads/master
  Home:   https://github.com/llvm/llvm-project
  Commit: d256797c9035aebf0309489c04dc34f8bae49dc4
      https://github.com/llvm/llvm-project/commit/d256797c9035aebf0309489c04dc34f8bae49dc4
  Author: JonChesterfield <jonathanchesterfield at gmail.com>
  Date:   2020-09-29 (Tue, 29 Sep 2020)

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

  Log Message:
  -----------
  [nfc][libomptarget] Drop parameter to named_sync

[nfc][libomptarget] Drop parameter to named_sync

named_sync has one call site (in sync.cu) where it always passed L1_BARRIER.
Folding this into the call site and dropping the macro is a simplification.

amdgpu doesn't have ptx' bar.sync instruction. A correct implementation of
__kmpc_impl_named_sync in terms of shared memory is much easier if it can
assume that the barrier argument is this constant. Said implementation is left
for a second patch.

Reviewed By: jdoerfert

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




More information about the All-commits mailing list