Skip to content

Commit 52315f9

Browse files
ddpaganddpagan
authored
[clang][OpenMP] Fix target data if/logical expression assert fail (#70268)
Fixed assertion failure Basic Block in function 'main' does not have terminator! label %land.end caused by premature setting of CodeGenIP upon entry to emitTargetDataCalls, where subsequent evaluation of logical expression created new basic blocks, leaving CodeGenIP pointing to the wrong basic block. CodeGenIP is now set near the end of the function, just prior to generating a comparison of the logical expression result (from the if clause) which uses CodeGenIP to insert new IR.
1 parent f24c443 commit 52315f9

File tree

2 files changed

+125
-5
lines changed

2 files changed

+125
-5
lines changed

‎clang/lib/CodeGen/CGOpenMPRuntime.cpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -10230,11 +10230,6 @@ void CGOpenMPRuntime::emitTargetDataCalls(
1023010230
PrePostActionTy NoPrivAction;
1023110231

1023210232
using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
10233-
InsertPointTy AllocaIP(CGF.AllocaInsertPt->getParent(),
10234-
CGF.AllocaInsertPt->getIterator());
10235-
InsertPointTy CodeGenIP(CGF.Builder.GetInsertBlock(),
10236-
CGF.Builder.GetInsertPoint());
10237-
llvm::OpenMPIRBuilder::LocationDescription OmpLoc(CodeGenIP);
1023810233

1023910234
llvm::Value *IfCondVal = nullptr;
1024010235
if (IfCond)
@@ -10314,6 +10309,11 @@ void CGOpenMPRuntime::emitTargetDataCalls(
1031410309
// Source location for the ident struct
1031510310
llvm::Value *RTLoc = emitUpdateLocation(CGF, D.getBeginLoc());
1031610311

10312+
InsertPointTy AllocaIP(CGF.AllocaInsertPt->getParent(),
10313+
CGF.AllocaInsertPt->getIterator());
10314+
InsertPointTy CodeGenIP(CGF.Builder.GetInsertBlock(),
10315+
CGF.Builder.GetInsertPoint());
10316+
llvm::OpenMPIRBuilder::LocationDescription OmpLoc(CodeGenIP);
1031710317
CGF.Builder.restoreIP(OMPBuilder.createTargetData(
1031810318
OmpLoc, AllocaIP, CodeGenIP, DeviceID, IfCondVal, Info, GenMapInfoCB,
1031910319
/*MapperFunc=*/nullptr, BodyCB, DeviceAddrCB, CustomMapperCB, RTLoc));

‎clang/test/OpenMP/target_data_if_logical_codegen.cpp

Lines changed: 120 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,120 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --version 3
2+
// REQUIRES: amdgpu-registered-target
3+
4+
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - \
5+
// RUN: | FileCheck %s
6+
7+
// Check same results after serialization round-trip
8+
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s
9+
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -emit-llvm %s -o - \
10+
// RUN: | FileCheck %s
11+
12+
// expected-no-diagnostics
13+
#ifndef HEADER
14+
#define HEADER
15+
16+
extern bool foo(bool);
17+
18+
int if_logical() {
19+
bool a = foo(true);
20+
bool b = foo(true);
21+
int pp = 42;
22+
int *p = &pp;
23+
#pragma omp target data if(a && b) map(to: p[0])
24+
{
25+
p[0]++;
26+
}
27+
if (p[0])
28+
return 1;
29+
return 0;
30+
}
31+
32+
int main() {
33+
return if_logical();
34+
}
35+
36+
#endif
37+
// CHECK-LABEL: define dso_local noundef i32 @_Z10if_logicalv(
38+
// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
39+
// CHECK-NEXT: entry:
40+
// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4
41+
// CHECK-NEXT: [[A:%.*]] = alloca i8, align 1
42+
// CHECK-NEXT: [[B:%.*]] = alloca i8, align 1
43+
// CHECK-NEXT: [[PP:%.*]] = alloca i32, align 4
44+
// CHECK-NEXT: [[P:%.*]] = alloca ptr, align 8
45+
// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
46+
// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
47+
// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
48+
// CHECK-NEXT: [[CALL:%.*]] = call noundef zeroext i1 @_Z3foob(i1 noundef zeroext true)
49+
// CHECK-NEXT: [[FROMBOOL:%.*]] = zext i1 [[CALL]] to i8
50+
// CHECK-NEXT: store i8 [[FROMBOOL]], ptr [[A]], align 1
51+
// CHECK-NEXT: [[CALL1:%.*]] = call noundef zeroext i1 @_Z3foob(i1 noundef zeroext true)
52+
// CHECK-NEXT: [[FROMBOOL2:%.*]] = zext i1 [[CALL1]] to i8
53+
// CHECK-NEXT: store i8 [[FROMBOOL2]], ptr [[B]], align 1
54+
// CHECK-NEXT: store i32 42, ptr [[PP]], align 4
55+
// CHECK-NEXT: store ptr [[PP]], ptr [[P]], align 8
56+
// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr [[A]], align 1
57+
// CHECK-NEXT: [[TOBOOL:%.*]] = trunc i8 [[TMP0]] to i1
58+
// CHECK-NEXT: br i1 [[TOBOOL]], label [[LAND_RHS:%.*]], label [[LAND_END:%.*]]
59+
// CHECK: land.rhs:
60+
// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr [[B]], align 1
61+
// CHECK-NEXT: [[TOBOOL3:%.*]] = trunc i8 [[TMP1]] to i1
62+
// CHECK-NEXT: br label [[LAND_END]]
63+
// CHECK: land.end:
64+
// CHECK-NEXT: [[TMP2:%.*]] = phi i1 [ false, [[ENTRY:%.*]] ], [ [[TOBOOL3]], [[LAND_RHS]] ]
65+
// CHECK-NEXT: br i1 [[TMP2]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_ELSE:%.*]]
66+
// CHECK: omp_if.then:
67+
// CHECK-NEXT: [[TMP3:%.*]] = load ptr, ptr [[P]], align 8
68+
// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[P]], align 8
69+
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP4]], i64 0
70+
// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
71+
// CHECK-NEXT: store ptr [[TMP3]], ptr [[TMP5]], align 8
72+
// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
73+
// CHECK-NEXT: store ptr [[ARRAYIDX]], ptr [[TMP6]], align 8
74+
// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
75+
// CHECK-NEXT: store ptr null, ptr [[TMP7]], align 8
76+
// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
77+
// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
78+
// CHECK-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 1, ptr [[TMP8]], ptr [[TMP9]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
79+
// CHECK-NEXT: br label [[OMP_IF_END:%.*]]
80+
// CHECK: omp_if.else:
81+
// CHECK-NEXT: br label [[OMP_IF_END]]
82+
// CHECK: omp_if.end:
83+
// CHECK-NEXT: [[TMP10:%.*]] = load ptr, ptr [[P]], align 8
84+
// CHECK-NEXT: [[ARRAYIDX4:%.*]] = getelementptr inbounds i32, ptr [[TMP10]], i64 0
85+
// CHECK-NEXT: [[TMP11:%.*]] = load i32, ptr [[ARRAYIDX4]], align 4
86+
// CHECK-NEXT: [[INC:%.*]] = add nsw i32 [[TMP11]], 1
87+
// CHECK-NEXT: store i32 [[INC]], ptr [[ARRAYIDX4]], align 4
88+
// CHECK-NEXT: br i1 [[TMP2]], label [[OMP_IF_THEN5:%.*]], label [[OMP_IF_ELSE6:%.*]]
89+
// CHECK: omp_if.then5:
90+
// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
91+
// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
92+
// CHECK-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP12]], ptr [[TMP13]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
93+
// CHECK-NEXT: br label [[OMP_IF_END7:%.*]]
94+
// CHECK: omp_if.else6:
95+
// CHECK-NEXT: br label [[OMP_IF_END7]]
96+
// CHECK: omp_if.end7:
97+
// CHECK-NEXT: [[TMP14:%.*]] = load ptr, ptr [[P]], align 8
98+
// CHECK-NEXT: [[ARRAYIDX8:%.*]] = getelementptr inbounds i32, ptr [[TMP14]], i64 0
99+
// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr [[ARRAYIDX8]], align 4
100+
// CHECK-NEXT: [[TOBOOL9:%.*]] = icmp ne i32 [[TMP15]], 0
101+
// CHECK-NEXT: br i1 [[TOBOOL9]], label [[IF_THEN:%.*]], label [[IF_END:%.*]]
102+
// CHECK: if.then:
103+
// CHECK-NEXT: store i32 1, ptr [[RETVAL]], align 4
104+
// CHECK-NEXT: br label [[RETURN:%.*]]
105+
// CHECK: if.end:
106+
// CHECK-NEXT: store i32 0, ptr [[RETVAL]], align 4
107+
// CHECK-NEXT: br label [[RETURN]]
108+
// CHECK: return:
109+
// CHECK-NEXT: [[TMP16:%.*]] = load i32, ptr [[RETVAL]], align 4
110+
// CHECK-NEXT: ret i32 [[TMP16]]
111+
//
112+
//
113+
// CHECK-LABEL: define dso_local noundef i32 @main(
114+
// CHECK-SAME: ) #[[ATTR3:[0-9]+]] {
115+
// CHECK-NEXT: entry:
116+
// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4
117+
// CHECK-NEXT: store i32 0, ptr [[RETVAL]], align 4
118+
// CHECK-NEXT: [[CALL:%.*]] = call noundef i32 @_Z10if_logicalv()
119+
// CHECK-NEXT: ret i32 [[CALL]]
120+
//

0 commit comments

Comments
 (0)