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

[Alexey Bataev via Phabricator via cfe-commits]

ABataev created this revision.
ABataev added a reviewer: jdoerfert.
Herald added subscribers: guansong, yaxunl.
ABataev requested review of this revision.
Herald added a subscriber: sstefan1.
Herald added a project: clang.

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.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D99445

Files:
  clang/lib/CodeGen/CGStmtOpenMP.cpp
  clang/test/OpenMP/nvptx_target_codegen.cpp


Index: clang/test/OpenMP/nvptx_target_codegen.cpp
===================================================================
--- clang/test/OpenMP/nvptx_target_codegen.cpp
+++ clang/test/OpenMP/nvptx_target_codegen.cpp
@@ -394,6 +394,12 @@
   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 @@
 // 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*,
Index: clang/lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -5742,6 +5742,7 @@
     CGF.CGM.getOpenMPRuntime().adjustTargetSpecificDataForLambdas(CGF, S);
 
   CGF.EmitStmt(S.getCapturedStmt(OMPD_target)->getCapturedStmt());
+  CGF.EnsureInsertPoint();
 }
 
 void CodeGenFunction::EmitOMPTargetDeviceFunction(CodeGenModule &CGM,


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D99445.333634.patch
Type: text/x-patch
Size: 1528 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210326/d355bff7/attachment-0001.bin>