From dcf96178cb3490028d9e2937cd8908b7c8fc75e6 Mon Sep 17 00:00:00 2001
From: Alexey Bataev <a.bataev@outlook.com>
Date: Fri, 26 Mar 2021 14:25:18 -0700
Subject: [PATCH] [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
---
 clang/lib/CodeGen/CGStmtOpenMP.cpp         |  1 +
 clang/test/OpenMP/nvptx_target_codegen.cpp | 12 ++++++++++++
 2 files changed, 13 insertions(+)

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*,