[Attributor]Fix PR51249: Stores from memory intrinsics incorrectly being replaced by undef

[jhuber6]

Summary

This patch fixes a bug observed in OpenMP reductions that occured when the
Attributor would attempt to simplify loads. The attributor would incorrectly
think that the memory intrinsic did not write to memory, and would then replace
the associated load with undef which would be turned to NaN in the final
output. This behaviour is observed in https://godbolt.org/z/j531848dv on line

314. This patch fixes https://bugs.llvm.org/show_bug.cgi?id=48607.

[ye-luo]

Please add a test to cover PR51249 case. Probably there are already tests for complex type but "-O3 -ffast-math" was not covered.

[jhuber6]

Adding test for bug.

[ye-luo]

LGTM

[ggeorgakoudis]

Great! Thanks Joseph.

[jhuber6]

I think there might be an underlying problem below this one. I'm going to continue looking through it.

[ye-luo]

In D107579#2929342, @jhuber6 wrote:

I think there might be an underlying problem below this one. I'm going to continue looking through it.

Does this patch have negative impact such that you are hesitating to land it and prefer investigating the underlying problem before landing?
If not, could you please land this patch?

[jhuber6]

In D107579#2935514, @ye-luo wrote:

In D107579#2929342, @jhuber6 wrote:

I think there might be an underlying problem below this one. I'm going to continue looking through it.

Does this patch have negative impact such that you are hesitating to land it and prefer investigating the underlying problem before landing?
If not, could you please land this patch?

Yes, I think this treats a secondary effect. I can land a patch that specifically ignores this case temporarily just to get it to work.

[lebedev.ri]

Please add an IR-level test for the attributor pass itself.

[jdoerfert]

Categorizing this as READ is correct. The error occurs without memset as well, see D107798.

[jhuber6]

Abandon in favor of D107798.

llvm/lib/Transforms/IPO/AttributorAttributes.cpp

if (auto *MI = dyn_cast_or_null<MemIntrinsic>(getCtxI())) {
         LengthVal = Length->getSExtValue();
       Value &Ptr = getAssociatedValue();
       unsigned ArgNo = getIRPosition().getCallSiteArgNo();
       ChangeStatus Changed;
       if (ArgNo == 0) {
         handleAccess(A, *MI, Ptr, nullptr, AccessKind::AK_WRITE, 0, Changed,
                      nullptr, LengthVal);
       } else if (ArgNo == 1) {
-        handleAccess(A, *MI, Ptr, nullptr, AccessKind::AK_READ, 0, Changed,
-                     nullptr, LengthVal);
+        handleAccess(A, *MI, Ptr, nullptr, AccessKind::AK_READ_WRITE, 0,
+                     Changed, nullptr, LengthVal);
       } else {
         LLVM_DEBUG(dbgs() << "[AAPointerInfo] Unhandled memory intrinsic "
                           << *MI << "\n");
         return indicatePessimisticFixpoint();
       }
       return Changed;
     }
 

openmp/libomptarget/test/offloading/bug51249.cpp

+// RUN: %libomptarget-compilexx-generic -O3 && %libomptarget-run-generic
+
+#include <complex>
+#include <cstdio>
+
+constexpr int N = 1024;
+
+std::complex<double> reduce(std::complex<double> *X, int N) {
+  std::complex<double> Sum = {0, 0};
+#pragma omp target teams distribute parallel for map(to: X[0:N]) reduction(+ : Sum)
+  for (int i = 0; i < N; ++i)
+    Sum += X[i];
+
+  return Sum;
+}
+
+int main() {
+  std::complex<double> X[N];
+
+  for (int i = 0; i < N; ++i)
+    X[i] = {1, -1};
+
+  std::complex<double> Sum = reduce(X, N);
+
+  // CHECK: PASS
+  if (std::real(Sum) + std::imag(Sum) == 0.0)
+    printf("PASS\n");
+}