[Openmp-commits] [openmp] e20d64c - [Clang][OpenMP] Fixed an issue of segment fault when using target nowait

[Shilei Tian via Openmp-commits]

Author: Shilei Tian
Date: 2020-10-26T22:33:22-04:00
New Revision: e20d64c3d9d81cad701f31d8481367222c76c787

URL: https://github.com/llvm/llvm-project/commit/e20d64c3d9d81cad701f31d8481367222c76c787
DIFF: https://github.com/llvm/llvm-project/commit/e20d64c3d9d81cad701f31d8481367222c76c787.diff

LOG: [Clang][OpenMP] Fixed an issue of segment fault when using target nowait

The implementation of target nowait just wraps the target region into a task. The essential four parameters (base ptr, ptr, size, mapper) are taken as firstprivate such that they will be copied to the private location. When there is no user-defined mapper, the mapper variable will be nullptr. However, it will be still copied to the corresponding place. Therefore, a memcpy will be generated and the source pointer will be nullptr, causing a segmentation fault. The root cause is when calling `emitOffloadingArraysArgument`, the last argument `Options` has a field about whether it requires a task. It only takes depend clause into account. In this patch, the nowait clause is also included.

There're two things that will be done in another patches:
1. target data nowait has not been supported yet. D90099 added the support.
2. When there is no mapper, the mapper array can be nullptr no matter whether it requires outer task or not. It can avoid an unnecessary data copy. This is an optimization that is covered in D90101.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D89844

Added: 
    openmp/libomptarget/test/offloading/bug47654.cpp

Modified: 
    clang/lib/CodeGen/CGOpenMPRuntime.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index cbeab66aa0f4..1e3937926ef0 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9664,11 +9664,10 @@ void CGOpenMPRuntime::emitTargetCall(
     TargetDataInfo Info;
     // Fill up the arrays and create the arguments.
     emitOffloadingArrays(CGF, CombinedInfo, Info);
-    bool HasDependClauses = D.hasClausesOfKind<OMPDependClause>();
     emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
                                  Info.PointersArray, Info.SizesArray,
                                  Info.MapTypesArray, Info.MappersArray, Info,
-                                 {/*ForEndTask=*/false, HasDependClauses});
+                                 {/*ForEndTask=*/false, RequiresOuterTask});
     InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
     InputInfo.BasePointersArray =
         Address(Info.BasePointersArray, CGM.getPointerAlign());

diff  --git a/openmp/libomptarget/test/offloading/bug47654.cpp b/openmp/libomptarget/test/offloading/bug47654.cpp
new file mode 100644
index 000000000000..14c8a9dde16b
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/bug47654.cpp
@@ -0,0 +1,29 @@
+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <cassert>
+#include <iostream>
+
+int main(int argc, char *argv[]) {
+  int i = 0, j = 0;
+
+#pragma omp target map(tofrom : i, j) nowait
+  {
+    i = 1;
+    j = 2;
+  }
+
+#pragma omp taskwait
+
+  assert(i == 1);
+  assert(j == 2);
+
+  std::cout << "PASS\n";
+
+  return 0;
+}
+
+// CHECK: PASS