[llvm-bugs] [Bug 49979] New: Offload Hierarchical Reduction on multiple Thread Give incorrect result

[via llvm-bugs]

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

            Bug ID: 49979
           Summary: Offload Hierarchical Reduction on multiple Thread 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

# Overview

Running an OpenMP reduction on the host where the body/kernel is a offload
hierarchical reduction (first reduction on team, then reduction on parallel)
produces a race condition / incorrect value.

## Note
- Work with `gcc`.
- Tested on NVIDIA hardware.
- Remove the first reduction give the correct value. 
- Merge the 2 offload reductions give the correct value.

# How to reproduce

```
cat test_parallel_teams_distribute__parallel_for.cpp
#include <iostream>
#include <cstdlib>
#include <cmath>
bool almost_equal(float x, float gold, float tol) {
  if ( std::signbit(x) != std::signbit(gold) )
  {
    x = std::abs(gold) - std::abs(x);
  }
  return std::abs(gold) * (1-tol) <= std::abs(x) && std::abs(x) <=
std::abs(gold) * (1 + tol);
}
void test_parallel_teams_distribute__parallel_for() {
  const int N0 { 32 };
  const int N1 { 32 };
  const int N2 { 32 };
  const float expected_value { N0*N1*N2 };
  float counter_N0{};
  #pragma omp parallel for reduction(+:counter_N0)
  for (int i0 = 0 ; i0 < N0 ; i0++ )
  {
    #pragma omp target teams distribute reduction(+: counter_N0)
    for (int i1 = 0 ; i1 < N1 ; i1++ )
    {
      #pragma omp parallel for reduction(+: counter_N0)
      for (int i2 = 0 ; i2 < N2 ; i2++ )
      {
          counter_N0 = counter_N0 + 1;
      }
    }
  }
  if (!almost_equal(counter_N0, expected_value, 0.1)) {
    std::cerr << "Expected: " << expected_value << " Got: " << counter_N0 <<
std::endl;
    std::exit(112);
  }
}
int main()
{
    test_parallel_teams_distribute__parallel_for();
}
> clang --version | head -n1
clang version 13.0.0 (https://github.com/llvm/llvm-project.git
ea14df695ebde3f5bdd5ba7548d3d49e8f1c4411)
> clang++ -fopenmp -fopenmp-targets=nvptx64 test_parallel_teams_distribute__parallel_for.cpp
> ./a.out
Expected: 32768 Got: 23552
> ./a.out
Expected: 32768 Got: 11264
```

-- 
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/20210415/68e80802/attachment.html>