[PATCH] D38978: [OpenMP] Enable the lowering of implicitly shared variables in OpenMP GPU-offloaded target regions to the GPU shared memory

[Gheorghe-Teodor Bercea via Phabricator via cfe-commits]

gtbercea created this revision.
Herald added subscribers: mgorny, jholewinski.

This patch is part of the development effort to add support in the current OpenMP GPU offloading implementation for implicitly sharing variables between a target region executed by the team master thread and the worker threads within that team.

This patch is the second of three required for successfully performing the implicit sharing of master thread variables with the worker threads within a team:
-Patch https://reviews.llvm.org/D38976 extends the CLANG code generation with code that handles shared variables.
-Patch (coming soon) extends the functionality of libomptarget to maintain a list of references to shared variables.

This patch adds a shared memory stack to the prolog of the kernel function representing the device offloaded OpenMP target region. The new passes along with the changes to existing ones, ensure that any OpenMP variable which needs to be shared across several threads will be allocated in this new stack, in the shared memory of the device. This patch covers the case of sharing variables from the master thread to the worker threads:

  #pragma omp target
  {
     // master thread only
     int v;
     #pragma omp parallel
     {
        // worker threads
        // use v
     }
  }


Repository:
  rL LLVM

https://reviews.llvm.org/D38978

Files:
  include/llvm/CodeGen/TargetPassConfig.h
  lib/CodeGen/TargetPassConfig.cpp
  lib/Target/NVPTX/CMakeLists.txt
  lib/Target/NVPTX/NVPTX.h
  lib/Target/NVPTX/NVPTXAsmPrinter.cpp
  lib/Target/NVPTX/NVPTXAssignValidGlobalNames.cpp
  lib/Target/NVPTX/NVPTXFrameLowering.cpp
  lib/Target/NVPTX/NVPTXFunctionDataSharing.cpp
  lib/Target/NVPTX/NVPTXFunctionDataSharing.h
  lib/Target/NVPTX/NVPTXInstrInfo.td
  lib/Target/NVPTX/NVPTXLowerAlloca.cpp
  lib/Target/NVPTX/NVPTXLowerSharedFrameIndicesPass.cpp
  lib/Target/NVPTX/NVPTXRegisterInfo.cpp
  lib/Target/NVPTX/NVPTXRegisterInfo.h
  lib/Target/NVPTX/NVPTXRegisterInfo.td
  lib/Target/NVPTX/NVPTXTargetMachine.cpp
  lib/Target/NVPTX/NVPTXUtilities.cpp
  lib/Target/NVPTX/NVPTXUtilities.h

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D38978.119210.patch
Type: text/x-patch
Size: 34686 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20171016/d9b38ffb/attachment-0001.bin>