This is an archive of the discontinued LLVM Phabricator instance.

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

Authored by cchen on Nov 3 2020, 11:50 AM.

Details

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

Diff Detail

Event Timeline

cchen created this revision.Nov 3 2020, 11:50 AM
Herald added a project: Restricted Project. · View Herald TranscriptNov 3 2020, 11:50 AM
cchen requested review of this revision.Nov 3 2020, 11:50 AM

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 updated this revision to Diff 302672.Nov 3 2020, 1:17 PM

Separate test to a independent file

cchen edited the summary of this revision. (Show Details)Nov 3 2020, 1:18 PM
cchen added a project: Restricted Project.
cchen edited the summary of this revision. (Show Details)
cchen added a comment.Nov 3 2020, 1:22 PM

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.

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

cchen added a comment.Nov 3 2020, 3:25 PM

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 added inline comments.Nov 4 2020, 7:24 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
2958–2970

Just:

if (hasTargetRegionEntryInfo(DeviceID, FileID, ParentName, LineNum))
  return;
ABataev added inline comments.Nov 4 2020, 7:55 AM
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 added inline comments.Nov 4 2020, 9:37 AM
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 added inline comments.Nov 4 2020, 9:45 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
2958–2970

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

ABataev added inline comments.Nov 4 2020, 9:47 AM
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 updated this revision to Diff 302882.Nov 4 2020, 10:01 AM
cchen edited the summary of this revision. (Show Details)

Fix based on feedback

cchen updated this revision to Diff 302884.Nov 4 2020, 10:03 AM

Fix coding style

ABataev added inline comments.Nov 4 2020, 10:12 AM
clang/lib/CodeGen/CGOpenMPRuntime.h
617 ↗(On Diff #302884)

IgnoreAddressId maybe?

cchen updated this revision to Diff 302898.Nov 4 2020, 10:30 AM

Refactor

This revision is now accepted and ready to land.Nov 4 2020, 10:31 AM
This revision was landed with ongoing or failed builds.Nov 4 2020, 10:37 AM
This revision was automatically updated to reflect the committed changes.