[llvm-bugs] [Bug 47039] New: Target followed by Atomic give incorrect result

[via llvm-bugs]

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

            Bug ID: 47039
           Summary: Target followed by Atomic give incorrect result
           Product: OpenMP
           Version: unspecified
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: enhancement
          Priority: P
         Component: Runtime Library
          Assignee: unassignedbugs at nondot.org
          Reporter: tapplencourt at anl.gov
                CC: llvm-bugs at lists.llvm.org

Created attachment 23826
  --> https://bugs.llvm.org/attachment.cgi?id=23826&action=edit
Reproducer

Found in `clang version 12.0.0 (https://github.com/llvm/llvm-project.git
55ead5bfffdc00e84cff347ee98471b5616a9f48)` running on a Ndvida V100

Hi,

When run, the following code:

```
  float counter_target{};
  #pragma omp target map(tofrom: counter_target)
  {
    #pragma omp atomic update
    counter_target = counter_target +  1. ;
  }
```
Produce 128, where we expected the result to be 1. Indeed only one thread
should be active in the target region.

LIBOMPTARGET_DEBUG=1 give us some hint from a potential root of the problem:
```
Target CUDA RTL --> Setting CUDA threads per block to default 128
Target CUDA RTL --> Using default number of teams 128
Target CUDA RTL --> Launch kernel with 128 blocks and 128 threads
```
It look like the runtime spwamned 128 threads, corresponding to 128 teams
regardless of the absence of the #teams pragma. 

Regards,
Thomas

PS: I joined a full reproducer for convenience. You can also found it here:
https://github.com/TApplencourt/OvO/blob/master/test_src/cpp/hierarchical_parallelism/atomic-float/target.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/20200807/7813cd2b/attachment-0001.html>