[PATCH] D93079: [OpenMP] Introduce an assumption to ignore possible external callers

[Johannes Doerfert via Phabricator via llvm-commits]

jdoerfert created this revision.
jdoerfert added reviewers: sstefan1, JonChesterfield, jhuber6, ggeorgakoudis.
Herald added subscribers: guansong, bollu, hiraditya, yaxunl.
jdoerfert requested review of this revision.
Herald added a project: LLVM.

In this patch we utilize the new assumption function attributes to
improve the OpenMPOpt behavior in the presence of externally visible
functions containing parallel regions when we compile for the GPU.

The existing OpenMP target state machine optimization will replace the
use of parallel region function pointers with dedicated handles if all
uses of the parallel region function pointer can be replaced. External
callers that are inside of OpenMP target regions will require the
function pointer to be used as there will be an indirect call. We now
allow the user to assume that external callers do not exist, or that
they might exist but not call from within an OpenMP (target) region.
The assumption in the IR is "omp_no_external_caller_in_target_region"
which can be created with OpenMP assumes

  omp begin assumes ext_omp_no_external_caller_in_target_region
  void foo() {
    #pragma omp parallel
    { ... }
  }
  omp end assumes

or the generic clang assume attribute

  void foo() __attribute__((assume("omp_no_external_caller_in_target_region"))) {
    #pragma omp parallel
    { ... }
  }


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D93079

Files:
  llvm/lib/Transforms/IPO/OpenMPOpt.cpp
  llvm/test/Transforms/OpenMP/gpu_state_machine_function_ptr_replacement.ll

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D93079.311069.patch
Type: text/x-patch
Size: 5158 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20201211/41ea41fc/attachment.bin>