[PATCH] D133694: [Clang][OpenMP] Fix use_device_addr

[Gheorghe-Teodor Bercea via Phabricator via cfe-commits]

doru1004 created this revision.
Herald added subscribers: guansong, yaxunl.
Herald added a project: All.
doru1004 requested review of this revision.
Herald added a reviewer: jdoerfert.
Herald added subscribers: cfe-commits, sstefan1.
Herald added a project: clang.

With the current implementation, the use_device_addr does not correctly use a pointer that was already mapped to the device and ends up in segmentation fault. The test showcases the situation.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D133694

Files:
  clang/lib/CodeGen/CGOpenMPRuntime.cpp
  clang/lib/CodeGen/CGStmtOpenMP.cpp
  clang/test/OpenMP/target_data_use_device_addr_codegen_ptr.cpp


Index: clang/test/OpenMP/target_data_use_device_addr_codegen_ptr.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/target_data_use_device_addr_codegen_ptr.cpp
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+int main() {
+    float x_array[256];
+    float *x = &x_array[0];
+
+    // make x available on the GPU
+    #pragma omp target data map(tofrom:x[0:256])
+    {
+        #pragma omp target data use_device_addr(x)
+        {
+            x[0] = 2;
+        }
+    }
+    return x[0] == 2;
+}
+
+// CHECK-LABEL: @main()
+// CHECK: [[X:%.+]] = alloca ptr, align 8
+// CHECK: call void @__tgt_target_data_begin_mapper(
+// CHECK: [[LOADED_X:%.+]] = load ptr, ptr [[X]], align 8
+// CHECK: [[BASE_PTR_GEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_BASE_PTR:%.+]], i32 0, i32 0
+// CHECK: store ptr [[LOADED_X]], ptr [[BASE_PTR_GEP]], align 8
+// CHECK: [[OFFLOAD_PTR_GEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_PTR:%.+]], i32 0, i32 0
+// CHECK: store ptr [[LOADED_X]], ptr [[OFFLOAD_PTR_GEP]], align 8
+// CHECK: getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_BASE_PTR]], i32 0, i32 0
+// CHECK: getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_PTR]], i32 0, i32 0
+// CHECK: call void @__tgt_target_data_begin_mapper(
+// CHECK: [[LOADED_DEVICE_X:%.+]] = load ptr, ptr [[BASE_PTR_GEP]], align 8
+// CHECK: %arrayidx5 = getelementptr inbounds float, ptr %13, i64 0
+// CHECK: store float 2.000000e+00, ptr %arrayidx5, align 4
+// CHECK: getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_BASE_PTR]], i32 0, i32 0
+// CHECK: getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_PTR]], i32 0, i32 0
+// CHECK: call void @__tgt_target_data_end_mapper(
+// CHECK: call void @__tgt_target_data_end_mapper(
+
+#endif
Index: clang/lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -7241,6 +7241,7 @@
     // declaration used by the mapping logic. In some cases we may get
     // OMPCapturedExprDecl that refers to the original declaration.
     const ValueDecl *MatchingVD = OrigVD;
+    bool isPartOfAStruct = false;
     if (const auto *OED = dyn_cast<OMPCapturedExprDecl>(MatchingVD)) {
       // OMPCapturedExprDecl are used to privative fields of the current
       // structure.
@@ -7248,6 +7249,7 @@
       assert(isa<CXXThisExpr>(ME->getBase()) &&
              "Base should be the current struct!");
       MatchingVD = ME->getMemberDecl();
+      isPartOfAStruct = true;
     }
 
     // If we don't have information about the current list item, move on to
@@ -7259,8 +7261,11 @@
     Address PrivAddr = InitAddrIt->getSecond();
     // For declrefs and variable length array need to load the pointer for
     // correct mapping, since the pointer to the data was passed to the runtime.
-    if (isa<DeclRefExpr>(Ref->IgnoreParenImpCasts()) ||
-        MatchingVD->getType()->isArrayType()) {
+    // Pointer types are already mapped correctly so no need to do a load unless
+    // the pointer type is part of a struct.
+    if ((isa<DeclRefExpr>(Ref->IgnoreParenImpCasts()) ||
+        MatchingVD->getType()->isArrayType()) &&
+        (isPartOfAStruct || !MatchingVD->getType()->isPointerType())) {
       QualType PtrTy = getContext().getPointerType(
           OrigVD->getType().getNonReferenceType());
       PrivAddr = EmitLoadOfPointer(
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8691,7 +8691,7 @@
           DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/true);
         } else {
           llvm::Value *Ptr;
-          if (IE->isGLValue())
+          if (IE->isGLValue() && !IE->getType()->isPointerType())
             Ptr = CGF.EmitLValue(IE).getPointer(CGF);
           else
             Ptr = CGF.EmitScalarExpr(IE);


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D133694.459439.patch
Type: text/x-patch
Size: 4235 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220912/9ad871ef/attachment.bin>