Author: Alexey Bataev Date: 2021-03-29T06:36:06-07:00 New Revision: dcf96178cb3490028d9e2937cd8908b7c8fc75e6
URL: https://github.com/llvm/llvm-project/commit/dcf96178cb3490028d9e2937cd8908b7c8fc75e6 DIFF: https://github.com/llvm/llvm-project/commit/dcf96178cb3490028d9e2937cd8908b7c8fc75e6.diff LOG: [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 Added: Modified: clang/lib/CodeGen/CGStmtOpenMP.cpp clang/test/OpenMP/nvptx_target_codegen.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index 2eaa481cd911..9a675a8870a2 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -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, diff --git a/clang/test/OpenMP/nvptx_target_codegen.cpp b/clang/test/OpenMP/nvptx_target_codegen.cpp index 56f04cb01f0a..7e32f7391f12 100644 --- a/clang/test/OpenMP/nvptx_target_codegen.cpp +++ b/clang/test/OpenMP/nvptx_target_codegen.cpp @@ -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*, _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
