maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[OPENMP]Fix PR49052: Clang crashed when compiling target code with assert(0). 

Need to insert a basic block during generation of the target region to
avoid crash for the GPU to be able always calling a cleanup action.
This cleanup action is required for the correct emission of the target
region for the GPU.

Differential Revision: https://reviews.llvm.org/D99445
This commit is contained in:
Alexey Bataev 2021-03-26 14:25:18 -07:00
parent 2f367f34fd
commit dcf96178cb
2 changed files with 13 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
clang/lib/CodeGen/CGStmtOpenMP.cpp
View File

@ -5742,6 +5742,7 @@ static void emitTargetRegion(CodeGenFunction &CGF, const OMPTargetDirective &S,
CGF.CGM.getOpenMPRuntime().adjustTargetSpecificDataForLambdas(CGF, S);
CGF.EmitStmt(S.getCapturedStmt(OMPD_target)->getCapturedStmt());
CGF.EnsureInsertPoint();
}
void CodeGenFunction::EmitOMPTargetDeviceFunction(CodeGenModule &CGM,

12
clang/test/OpenMP/nvptx_target_codegen.cpp
View File

@ -394,6 +394,12 @@ int baz(int f, double &a) {
return f;
}
extern void assert(int) throw() __attribute__((__noreturn__));
void unreachable_call() {
#pragma omp target
assert(0);
}
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+static.+347}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
@ -632,6 +638,12 @@ int baz(int f, double &a) {
// CHECK: [[RES:%.+]] = load i32, i32* [[RET]],
// CHECK: ret i32 [[RES]]
// CHECK: define {{.*}}void {{@__omp_offloading_.+unreachable_call.+l399}}()
// CHECK: call void @{{.*}}assert{{.*}}(i32 0)
// CHECK: unreachable
// CHECK: call void @__kmpc_kernel_deinit(i16 1)
// CHECK: ret void
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l331}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,