[OpenMP] target nested `use_device_ptr() if()` and is_device_ptr trigger asserts

[cchen]

Summary

Clang now asserts for the below case:

void clang::CodeGen::CGOpenMPRuntime::createOffloadEntriesAndInfoMetadata(): Assertion `std::get<0>(E) && "All ordered entries must exist!"' failed.

The reason why Clang hit the assert is because in
emitTargetDataCalls, both BeginThenGen and BeginElseGen call
registerTargetRegionEntryInfo and try to register the Entry in
OffloadEntriesTargetRegion with same key. If changing the expression in
if clause to any constant expression, then the assert disappear. (https://godbolt.org/z/TW7haj)

The assert itself is to avoid
user from accessing elements out of bound inside OrderedEntries in
createOffloadEntriesAndInfoMetadata.

In this patch, I add a check in registerTargetRegionEntryInfo to avoid
register the target region more than once.

A test case that triggers assert: https://godbolt.org/z/4cnGW8

[jdoerfert]

I do not understand the commit message. Can you try to make it clearer?

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

518

Create a new file please.

[cchen]

Separate test to a independent file

[cchen]

In D90704#2371938, @jdoerfert wrote:

I do not understand the commit message. Can you try to make it clearer?

Hi, sorry for the shabby commit message, I've updated it and use godbolt link instead of pasting the code snippet since the pragma is always removed automatically by git or arc.

[ABataev]

It would be good if you could identify the object which leads to a crash, I mean a target region, variable, etc.

[cchen]

In D90704#2372303, @ABataev wrote:

It would be good if you could identify the object which leads to a crash, I mean a target region, variable, etc.

 1	void
 2	add_one(float *b, int dm)
 3	{
 4	#pragma omp target data map(tofrom:b[:1]) use_device_ptr(b) if (dm)
 5	  {
 6	#pragma omp target is_device_ptr(b)
 7	  {
 8	    b[0] += 1;
 9	  }
10	  }
11	}

The code crashes at line 6. Below is the information I dumped from registerTargetRegionEntryInfo. (DeviceID, FileID, ParentName, LineNum, Addr, ID)

DeviceID: 16777220, FileID 38814071, ParentName add_one, LineNum: 6
Addr:
; Function Attrs: noinline norecurse nounwind optnone ssp uwtable
define internal void @__omp_offloading_1000004_2504177_add_one_l6(float* %b) #2 {
entry:
  %b.addr = alloca float*, align 8
  store float* %b, float** %b.addr, align 8
  %0 = load float*, float** %b.addr, align 8
  %arrayidx = getelementptr inbounds float, float* %0, i64 0
  %1 = load float, float* %arrayidx, align 4
  %add = fadd float %1, 1.000000e+00
  store float %add, float* %arrayidx, align 4
  ret void
}

ID:
@.__omp_offloading_1000004_2504177_add_one_l6.region_id = weak constant i8 0

I found that both the BeginThenGen lambda and BeginElseGen lambda calls registerTargetRegionEntryInfo and both do

OffloadEntryInfoTargetRegion Entry(OffloadingEntriesNum, Addr, ID, Flags);
OffloadEntriesTargetRegion[DeviceID][FileID][ParentName][LineNum] = Entry;
OffloadingEntriesNum++;

Which will lead to out-of-bounds access in createOffloadEntriesAndInfoMetadata since OffloadingEntriessNum is increased even if the size of OffloadEntriesTargetRegion is not increased.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

2958–2970

Just:

if (hasTargetRegionEntryInfo(DeviceID, FileID, ParentName, LineNum))
  return;

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

2958–2970

Actually, even better to do something like this:

if (Flags == OffloadEntriesInfoManagerTy::OMPTargetRegionEntryTargetRegion && hasTargetRegionEntryInfo(DeviceID, FileID, ParentName, LineNum))
  return;
assert(!hasTargetRegionEntryInfo(DeviceID, FileID, ParentName, LineNum) && "Target region entry already registered!");

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

2958–2970

I tried this code but found that hasTargetRegionEntryInfo is not doing what we want since we need the below code to return true in hasTargetRegionEntryInfo.

if (PerLine->second.getAddress() || PerLine->second.getID()) {
  return false;
}

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

2958–2970

Do we want to create a new helper function like hasTargetRegionEntryInfo or something?

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

2958–2970

Better to extend the old one, I think. Add a new parameter to check if we need to check for address and id

[cchen]

Fix based on feedback

[cchen]

Fix coding style

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.h

617 ↗

(On Diff #302884)

IgnoreAddressId maybe?

[cchen]

Refactor

[ABataev]

LG

clang/lib/CodeGen/CGOpenMPRuntime.cpp

if (CGM.getLangOpts().OpenMPIsDevice) {
     }
     auto &Entry =
         OffloadEntriesTargetRegion[DeviceID][FileID][ParentName][LineNum];
     assert(Entry.isValid() && "Entry not initialized!");
     Entry.setAddress(Addr);
     Entry.setID(ID);
     Entry.setFlags(Flags);
   } else {
+    auto PerDevice = OffloadEntriesTargetRegion.find(DeviceID);
+    if (PerDevice != OffloadEntriesTargetRegion.end()) {
+      auto PerFile = PerDevice->second.find(FileID);
+      if (PerFile != PerDevice->second.end()) {
+        auto PerParentName = PerFile->second.find(ParentName);
+        if (PerParentName != PerFile->second.end()) {
+          auto PerLine = PerParentName->second.find(LineNum);
+          if (PerLine != PerParentName->second.end()) {
+            return;
+          }
+        }
+      }
+    }
     OffloadEntryInfoTargetRegion Entry(OffloadingEntriesNum, Addr, ID, Flags);
     OffloadEntriesTargetRegion[DeviceID][FileID][ParentName][LineNum] = Entry;
     ++OffloadingEntriesNum;
   }
 }
 
 bool CGOpenMPRuntime::OffloadEntriesInfoManagerTy::hasTargetRegionEntryInfo(
     unsigned DeviceID, unsigned FileID, StringRef ParentName,

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

 };
 
 void bar(double *arg){
   ST<double> A(arg);
   A.foo(arg);
   ++arg;
 }
 #endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK3 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
+// RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-64
+// RUN: %clang_cc1 -DCK3 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-32
+// RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-32
+
+// RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s
+// RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
+// RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s
+// RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
+// SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
+#ifdef CK3
+
+// CK3: [[MTYPE00:@.+]] = {{.*}}constant [1 x i64] [i64 99]
+// CK3: [[MTYPE01:@.+]] = {{.*}}constant [1 x i64] [i64 288]
+// CK3: [[MTYPE02:@.+]] = {{.*}}constant [1 x i64] [i64 288]
+
+void add_one(float *b, int dm)
+{
+  // CK3:     [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK3:     [[CBP:%.+]] = bitcast i8** [[BP]] to float**
+  // CK3:     store float* [[B_ADDR:%.+]], float** [[CBP]]
+  // CK3:     call void @__tgt_target_data_begin{{.+}}[[MTYPE00]]
+  // CK3:     [[VAL:%.+]] = load float*, float** [[CBP]],
+  // CK3-NOT: store float* [[VAL]], float** [[DECL]],
+  // CK3:     store float* [[VAL]], float** [[PVT:%.+]],
+  // CK3:     [[TT:%.+]] = load float*, float** [[PVT]],
+  // CK3:     call i32 @__tgt_target{{.+}}[[MTYPE01]]
+  // CK3:     call i32 @__tgt_target{{.+}}[[MTYPE02]]
+  // CK3:     call void @__tgt_target_data_end{{.+}}[[MTYPE00]]
+#pragma omp target data map(tofrom:b[:1]) use_device_ptr(b) if(dm == 0)
+  {
+#pragma omp target is_device_ptr(b)
+  {
+    b[0] += 1;
+  }
+  }
+}
+
+#endif
 #endif