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

[ABataev]

Summary

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.

[jdoerfert]

LG

clang/lib/CodeGen/CGStmtOpenMP.cpp

static void emitTargetRegion(CodeGenFunction &CGF, const OMPTargetDirective &S,
   CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
   (void)CGF.EmitOMPFirstprivateClause(S, PrivateScope);
   CGF.EmitOMPPrivateClause(S, PrivateScope);
   (void)PrivateScope.Privatize();
   if (isOpenMPTargetExecutionDirective(S.getDirectiveKind()))
     CGF.CGM.getOpenMPRuntime().adjustTargetSpecificDataForLambdas(CGF, S);
 
   CGF.EmitStmt(S.getCapturedStmt(OMPD_target)->getCapturedStmt());
+  CGF.EnsureInsertPoint();
 }
 
 void CodeGenFunction::EmitOMPTargetDeviceFunction(CodeGenModule &CGM,
                                                   StringRef ParentName,
                                                   const OMPTargetDirective &S) {
   auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
     emitTargetRegion(CGF, S, Action);
   };

clang/test/OpenMP/nvptx_target_codegen.cpp

 }
 
 int baz(int f, double &a) {
 #pragma omp parallel
   f = 2 + 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*,
 // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
 // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
 //
 // CHECK: [[AWAIT_WORK]]

 // CHECK: store i32 [[RES]], i32* [[RET:%.+]],
 // CHECK: br i1 [[IS_SPMD]], label
 // CHECK: [[BC:%.+]] = bitcast [[GLOBAL_ST]]* [[ITEMS]] to i8*
 // CHECK: call void @__kmpc_data_sharing_pop_stack(i8* [[BC]])
 // CHECK: br label
 // 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*,
 // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
 // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
 // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
 //
 // CHECK: [[AWAIT_WORK]]