[llvm-bugs] [Bug 48851] New: Clang Incorrectly Determines SPMD vs Generic OpenMP Kernel

[via llvm-bugs]

https://bugs.llvm.org/show_bug.cgi?id=48851

            Bug ID: 48851
           Summary: Clang Incorrectly Determines SPMD vs Generic OpenMP
                    Kernel
           Product: clang
           Version: trunk
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: normal
          Priority: P
         Component: OpenMP
          Assignee: unassignedclangbugs at nondot.org
          Reporter: estewart08 at gmail.com
                CC: llvm-bugs at lists.llvm.org

Created attachment 24410
  --> https://bugs.llvm.org/attachment.cgi?id=24410&action=edit
SPMD vs Generic Bug Reproducer

According to the docs, https://clang.llvm.org/docs/OpenMPSupport.html:

The compiler will always attempt to use the SPMD mode wherever possible. SPMD
mode will not be used if: 
  The target region contains user code (other than OpenMP-specific 
  directives) in between the target and the parallel directives.

  #pragma omp target teams thread_limit(THREADS) num_teams(TEAMS)
map(from:gpu_results)
  {
    int dist[THREADS];
    // Uncomment line below to trigger generic kernel
    //dist[0] = 0;
    #pragma omp parallel
    {

Clang determines this is a SPMD kernel, even though there is user code between
the target and parallel pragmas, and the dist array becomes thread private.  In
the attached code, there is a reduction later on that assumes dist is team
private.  

When looking for the child of the target teams pragma, the compiler skips the
declaration of the dist array and sees the parallel pragma for SPMD mode.  This
occurs due to the logic in clang/lib/CodeGen/CGOpenMPRuntime.cpp on line 6525,
VD->getType().isTrivialType(Ctx). Since this declaration is of trivial type it
is skipped, which has adverse effects.  

To summarize:
SPMD (dist array is thread private)
int dist[THREADS]; //skipped
#pragma omp parallel //child that is returned

Generic (dist array is team private)
int dist[THREADS]; //skipped
dist[0] = 0 // child 1
#pragma omp parallel //child 2, multiple children, return nullptr

Attachment is derived from:
https://github.com/zjin-lcf/oneAPI-DirectProgramming/blob/master/all-pairs-distance-omp/main.cpp

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20210122/3f127728/attachment.html>