This is an archive of the discontinued LLVM Phabricator instance.

[OpenMP] Add implicit data sharing support when offloading to NVIDIA GPUs using OpenMP device offloading
ClosedPublic

Authored by gtbercea on Oct 16 2017, 2:21 PM.

Details

Summary

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:

  • Patch D38978 to the LLVM NVPTX backend which ensures the lowering of shared variables to an device memory which allows the sharing of references;
  • Patch (coming soon) is a patch to libomptarget runtime library 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.

Event Timeline

gtbercea created this revision.Oct 16 2017, 2:21 PM
This revision is now accepted and ready to land.Nov 3 2017, 12:38 PM
Hahnfeld added inline comments.Nov 3 2017, 12:48 PM
lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
2393

Is this meant to be a comment or should there be something happening here?

2396

Especially, NextBB can never be anything else than nullptr...

gtbercea updated this revision to Diff 121538.Nov 3 2017, 1:24 PM
gtbercea marked 2 inline comments as done.
gtbercea updated this revision to Diff 121543.Nov 3 2017, 1:48 PM
gtbercea edited the summary of this revision. (Show Details)

Remove blocks.

ABataev accepted this revision.Nov 20 2017, 6:36 AM

LG

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
474

Later we should remove it from the code.

gtbercea closed this revision.Nov 21 2017, 7:55 AM