[HeterogeneousDWARF] Workarounds to prevent CGDebugInfo crashes

Any path which can produce a DIExpression is invalid when
-gheterogeneous-dwarf=diexpr, so stub out these paths to avoid a crash
until they are implemented.

Change-Id: I81446f887727fb343a2a0527f50aa50f22978d81

Author: slinder1

clang/lib/CodeGen/CGDebugInfo.cpp

@@ -5515,6 +5515,9 @@ llvm::DIType *CGDebugInfo::CreateSelfType(const QualType &QualTy,
 void CGDebugInfo::EmitDeclareOfBlockDeclRefVariable(
     const VarDecl *VD, llvm::Value *Storage, CGBuilderTy &Builder,
     const CGBlockInfo &blockInfo, llvm::Instruction *InsertPoint) {
+  // FIXME: Workaround to prevent crash when using with -gheterogeneous-dwarf
+  if (CGM.getCodeGenOpts().isHeterogeneousDwarfEnabled())
+    return;
   assert(CGM.getCodeGenOpts().hasReducedDebugInfo());
   assert(!LexicalBlockStack.empty() && "Region stack mismatch, stack empty!");
 
@@ -5647,6 +5650,9 @@ void CGDebugInfo::EmitDeclareOfBlockLiteralArgVariable(const CGBlockInfo &block,
                                                        unsigned ArgNo,
                                                        llvm::AllocaInst *Alloca,
                                                        CGBuilderTy &Builder) {
+  // FIXME: Workaround to prevent crash when using with -gheterogeneous-dwarf
+  if (CGM.getCodeGenOpts().isHeterogeneousDwarfEnabled())
+    return;
   assert(CGM.getCodeGenOpts().hasReducedDebugInfo());
   ASTContext &C = CGM.getContext();
   const BlockDecl *blockDecl = block.getBlockDecl();
@@ -6458,6 +6464,12 @@ void CGDebugInfo::EmitGlobalVariableForHeterogeneousDwarf(
 
 void CGDebugInfo::EmitExternalVariable(llvm::GlobalVariable *Var,
                                        const VarDecl *D) {
+  // FIXME: Workaround to prevent crash when using with -gheterogeneous-dwarf
+  // NOTE: Only currently reachable for BPF target, but check added for
+  // completeness and in case this changes.
+  if (CGM.getCodeGenOpts().isHeterogeneousDwarfEnabled())
+    return;
+
   assert(CGM.getCodeGenOpts().hasReducedDebugInfo());
   if (D->hasAttr<NoDebugAttr>())
     return;

clang/test/CodeGen/debug-info-block-expr-heterogeneous-dwarf.c

@@ -35,7 +35,7 @@ void test_local_block() {
 
 // FIXME(KZHURAVL): Update EmitDeclareOfBlockDeclRefVariable and EmitDeclareOfBlockLiteralArgVariable.
 // CHECK-LABEL: @__test_local_block_block_invoke
-// CHECK: #dbg_declare({{.*}}!DIExpression(DW_OP_deref, DW_OP_plus_uconst, {{[0-9]+}}, DW_OP_deref, DW_OP_plus_uconst, {{[0-9]+}}, DW_OP_deref, DW_OP_plus_uconst, {{[0-9]+}}){{.*}})
+// CHECK-NOT: #dbg_declare({{.*}}
   ^ { block_var = 1; }();
 }
 

clang/test/CodeGenOpenCL/amdgpu-debug-info-blocks-heterogeneous-dwarf.cl

@@ -0,0 +1,14 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang -Xclang -cl-std=CL2.0 -emit-llvm -g -gheterogeneous-dwarf=diexpr -O0 -S -nogpulib -target amdgcn-amd-amdhsa -mcpu=gfx1010 -o - %s | FileCheck %s
+// RUN: %clang -Xclang -cl-std=CL2.0 -emit-llvm -g -gheterogeneous-dwarf=diexpr -O0 -S -nogpulib -target amdgcn-amd-amdhsa-opencl -mcpu=gfx1010 -o - %s | FileCheck %s
+
+// FIXME: Currently just testing that we don't crash; test for the absense
+// of meaningful debug information for the block is to identify this test
+// to update/replace when this is implemented.
+
+// CHECK-NOT: call void @llvm.dbg.{{.*}}%block.capture.addr
+
+void fn(__global uint* res);
+__kernel void kern(__global uint* res) {
+  void (^kernelBlock)(void) = ^{ fn(res); };
+}

clang/test/CodeGenOpenCL/bpf-debug-info-extern-heterogeneous-dwarf.cl

@@ -0,0 +1,13 @@
+// REQUIRES: bpf-registered-target
+// RUN: %clang -Xclang -cl-std=CL2.0 -emit-llvm -g -gheterogeneous-dwarf=diexpr -O0 -S -nogpulib -target bpf-linux-gnu -o - %s | FileCheck %s
+
+// FIXME: Currently just testing that we don't crash; test for the absense
+// of meaningful debug information for the extern is to identify this test
+// to update/replace when this is implemented.
+
+// CHECK-NOT: DIGlobalVariable
+
+extern char ch;
+int test() {
+  return ch;
+}