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

[Chi Chun Chen via Phabricator via cfe-commits]

cchen added a comment.

In D90704#2372303 <https://reviews.llvm.org/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.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D90704/new/

https://reviews.llvm.org/D90704