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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
