[PATCH] D38976: [OpenMP] Add implicit data sharing support when offloading to NVIDIA GPUs using OpenMP device offloading

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

gtbercea created this revision.
Herald added a subscriber: 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 first of three required for successfully performing the implicit sharing of master thread variables with the worker threads within a team. The remaining two patches are:

- a patch to the LLVM NVPTX backend which ensures the lowering of shared variables to an device memory which allows the sharing of references;
- a runtime patch to libomptarget which ensures that a list of references to shared variables is properly maintained.

A simple code snippet which illustrates an implicit data sharing situation is as follows:

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

Variable v is implicitly shared from the team master thread which executes the code in between the target and parallel directives. The worker threads must operate on the latest version of v, including any updates performed by the master.

The code generated in this patch relies on the LLVM NVPTX patch (mentioned above) which prevents v from being lowered in the thread local memory of the master thread thus making the reference to this variable un-shareable with the workers. This ensures that the code generated by this patch is correct.
Since the parallel region is outlined the passing of arguments to the outlined regions must preserve the original order of arguments. The runtime therefore maintains a list of references to shared variables thus ensuring their passing in the correct order. The passing of arguments to the outlined parallel function is performed in a separate function which the data sharing infrastructure constructs in this patch. The function is inlined when optimizations are enabled.


Repository:
  rL LLVM

https://reviews.llvm.org/D38976

Files:
  lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
  lib/CodeGen/CGOpenMPRuntimeNVPTX.h
  test/OpenMP/nvptx_data_sharing.cpp
  test/OpenMP/nvptx_parallel_codegen.cpp
  test/OpenMP/nvptx_target_teams_codegen.cpp

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D38976.119211.patch
Type: text/x-patch
Size: 20696 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20171016/6fd0b2a5/attachment-0001.bin>