Squashed version of diexpression-poison patches

This is a combination of multiple commits.

This is the 1st commit message:

DIOp-based DIExpression infrastructure

Add the minimal support for DIOp-in-DIExpression, including DIOpFragment
and DW_OP_LLVM_poisoned, and the API to differentiate the variants.

This is the commit message #2:

[Bitcode] Serialization for DIOp-based DIExpression

This is the commit message #3:

Dwarf generation for DIOp-based DIExpression

This is a rather large patch with very minimal testing. It should
probably be split up further, and more tests exercising every path are
needed.

This is the commit message #4:

Extend clang option and add clang codegen for DIOp-based DIExpression

This is the commit message #5:

Support DIOp-based DIExpressions in SROA/mem2reg/instcombine

This patch doesn't actually touch those passes, but just the utilities
they use, namely createFragmentExpression(), and
ConvertDebugDeclareToDebugValue().

This doesn't include assignment tracking, which has some special handling in
SROA.cpp and PromoteMemoryToRegister.cpp. We're not planning on using dbg.assign
for this (at least for the time being), so I just ignored that for now.

This is the commit message #6:

[HeterogeneousDwarf] Handle signed integers in DIOpShr and DIOpConvert

This commit adds new DI operations to differentiate between a zext/sext
DIOpConvert and a ashr/lshr DIOpShr. It isn't possible to use the IR
type for this, since it doesn't distinguish between signed and unsigned
integers. Fixes SWDEV-466183.

This is the commit message #7:

[Debugify] Add a flag to make DIOp-based DIExpressions

This should be useful for porting debugify-based optimizer tests. Part
of SWDEV-462843.

This is the commit message #8:

Handle new DIOp-DIExpressions in replaceAllDbgUsesWith

Fixes part of SWDEV-465029

This is the commit message #9:

Handle new DIOp-DIExpressions in salvageDebugInfo

This fixes part of SWDEV-465029.

This is the commit message #10:

[IRGen] Strip addrspacecasts when creating dbg.declares

dbg.def does this in DIBuilder, but this commit just adds it to clang to
avoid introducing a diff with upstream.

This is the commit message #11:

Verifier support for DIOp-based DIExpression

Effectively a ported and updated version of
https://gerrit-git.amd.com/c/lightning/ec/llvm-project/+/974933

Changed to one overload set rather than distinct method names for
visitor base so the derived class can opt in to non-exhaustive visiting,
rather than it be implied. Added a means to visit the result of the
expression when it is otherwise valid (i.e. there is exactly one result).

Moved as much of the validation as possible into the base class, leaving
the only derived class using the visitor so far to essentially just do
bitsize-based type checks when the arguments and/or DataLayout are available.

The AsmPrinter support could be ported over to the visitor pattern
eventually, and the verifier can be ported over to DIExpr, but these
are left as future improvements.

This is the commit message #12:

Add DIOp AsmPrinter support for Convert/ZExt/SExt

Since AsmPrinter currently require values on evaluation stack to be of
generic type, we have to use the "legacy" dwarf-4 conversion operations.
This can be a little verbose, particularly for sext. It would be
technically possible to represent these with three DW_OP_converts
(converting generic -> signed FromBits -> signed ToBits -> generic), but
using the legacy version seemed simpler.

In the future we could use DW_OP_convert to implement these, but in
order to do that we would need to ensure that values on the dwarf
evaluation stack have non-generic types. For instance, we would need to
use use DW_OP_const_type instead of DW_OP_lit for constants. Failing to
do so would break binary operators, which require compatible types for
their inputs.

One note: it seems like it's ambigious whether a DIOpArg that produces a
negative signed value with a type smaller than the generic type will
have it's higher order bits signed extended or not. For constants,
FastISel produces a zero extended value, and non-fast ISel produces a
sign extended value (see FastISel.cpp:1263 vs InstrEmitter:740 @ this
commit). This can be observed by passing --fast-isel=false to the test
file. SExt is correct for both cases, and always creates a fully
sign-extended value of the generic type.

Fixes SWDEV-467965

This is the commit message #13:

Add DIOp-in-DIExpression test for MIR serialization

This is the commit message #14:

Change -gheterogeneous-dwarf default to diexpression

This is the commit message #15:

[HeterogeneousDWARF] Various fixes against PSDB

Resolve failures in PSDB smoke tests, catch2 tests, and one lit test
(caused by upstream work in SROA).

Several `FIXME(diexpression-poison)` comments mark places where there
is additional work required still, e.g. workarounds or partial fixes
to get changes passing PSDB.

This is the commit message #16:

[HeterogeneousDWARF] Restore -gheterogeneous-dwarf cc1 option

This is the commit message #17:

[MIR] Replace bespoke DIExpression parser

Resolve FIXME by using the LLParser implementation of parseDIExpression
from the MIParser.

This is the commit message #18:

[HetereogeneousDWARF] Revert default to =diexpr

Change-Id: I650ec1e9f6f88ef881f79ef3959785439871e0ba

Author: slinder1

clang/include/clang/Basic/CodeGenOptions.def

@@ -30,8 +30,6 @@ CODEGENOPT(Name, Bits, Default)
 CODEGENOPT(DisableIntegratedAS, 1, 0) ///< -no-integrated-as
 CODEGENOPT(RelaxELFRelocations, 1, 1) ///< -Wa,-mrelax-relocations={yes,no}
 CODEGENOPT(AsmVerbose        , 1, 0) ///< -dA, -fverbose-asm.
-CODEGENOPT(HeterogeneousDwarf, 1, 0) ///< Enable DWARF extensions for
-                                     ///< heterogeneous debugging.
 CODEGENOPT(PreserveAsmComments, 1, 1) ///< -dA, -fno-preserve-as-comments.
 CODEGENOPT(AssumeSaneOperatorNew , 1, 1) ///< implicit __attribute__((malloc)) operator new
 CODEGENOPT(AssumeUniqueVTables , 1, 1) ///< Assume a class has only one vtable.

clang/include/clang/Basic/CodeGenOptions.h

@@ -158,6 +158,21 @@ class CodeGenOptions : public CodeGenOptionsBase {
     Never,    // No loop is assumed to be finite.
   };
 
+  enum class HeterogeneousDwarfOpts {
+    Disabled,     //< Do not emit any heterogeneous dwarf metadata.
+    DIExpr,       //< Enable DIExpr-based metadata.
+    DIExpression, //< Enable DIExpression-based metadata.
+  };
+  bool isHeterogeneousDwarfEnabled() const {
+    return getHeterogeneousDwarfMode() != HeterogeneousDwarfOpts::Disabled;
+  }
+  bool isHeterogeneousDwarfDIExpr() const {
+    return getHeterogeneousDwarfMode() == HeterogeneousDwarfOpts::DIExpr;
+  }
+  bool isHeterogeneousDwarfDIExpression() const {
+    return getHeterogeneousDwarfMode() == HeterogeneousDwarfOpts::DIExpression;
+  }
+
   enum AssignmentTrackingOpts {
     Disabled,
     Enabled,

clang/include/clang/Basic/DebugOptions.def

@@ -69,6 +69,10 @@ BENIGN_DEBUGOPT(NoInlineLineTables, 1, 0) ///< Whether debug info should contain
 
 DEBUGOPT(DebugStrictDwarf, 1, 1) ///< Whether or not to use strict DWARF info.
 
+/// Control DWARF extensions for heterogeneous debugging enablement and approach.
+BENIGN_ENUM_DEBUGOPT(HeterogeneousDwarfMode, HeterogeneousDwarfOpts, 2,
+                     HeterogeneousDwarfOpts::Disabled)
+
 /// Control the Assignment Tracking debug info feature.
 BENIGN_ENUM_DEBUGOPT(AssignmentTrackingMode, AssignmentTrackingOpts, 2,
                      AssignmentTrackingOpts::Disabled)

clang/include/clang/Driver/Options.td

@@ -4211,12 +4211,23 @@ def gdwarf64 : Flag<["-"], "gdwarf64">, Group<g_Group>,
 def gdwarf32 : Flag<["-"], "gdwarf32">, Group<g_Group>,
   Visibility<[ClangOption, CC1Option, CC1AsOption]>,
   HelpText<"Enables DWARF32 format for ELF binaries, if debug information emission is enabled.">;
-def gheterogeneous_dwarf : Flag<["-"], "gheterogeneous-dwarf">,
+
+def gheterogeneous_dwarf_EQ : Joined<["-"], "gheterogeneous-dwarf=">,
   Group<g_Group>, Visibility<[ClangOption, CC1Option]>,
-  HelpText<"Enable DWARF extensions for heterogeneous debugging">,
-  MarshallingInfoFlag<CodeGenOpts<"HeterogeneousDwarf">>;
+  HelpText<"Control DWARF extensions for heterogeneous debugging">,
+  Values<"disabled,diexpr,diexpression">,
+  NormalizedValuesScope<"CodeGenOptions::HeterogeneousDwarfOpts">,
+  NormalizedValues<["Disabled","DIExpr","DIExpression"]>,
+  MarshallingInfoEnum<CodeGenOpts<"HeterogeneousDwarfMode">, "Disabled">;
+def gheterogeneous_dwarf : Flag<["-"], "gheterogeneous-dwarf">, Group<g_Group>,
+  Visibility<[ClangOption, CC1Option]>,
+  HelpText<"Enable DIExpr-based DWARF extensions for heterogeneous debugging">,
+  Alias<gheterogeneous_dwarf_EQ>, AliasArgs<["diexpr"]>;
 def gno_heterogeneous_dwarf : Flag<["-"], "gno-heterogeneous-dwarf">,
-  Group<g_Group>, HelpText<"Disable DWARF extensions for heterogeneous debugging">;
+  Visibility<[ClangOption, CC1Option]>,
+  Group<g_Group>,
+  HelpText<"Disable DWARF extensions for heterogeneous debugging">,
+  Alias<gheterogeneous_dwarf_EQ>, AliasArgs<["disabled"]>;
 
 def gcodeview : Flag<["-"], "gcodeview">,
   HelpText<"Generate CodeView debug information">,

clang/lib/CodeGen/CGDebugInfo.cpp

@@ -777,7 +777,20 @@ class CGDebugInfo {
   /// anonymous decl and create static variables for them. The first
   /// time this is called it needs to be on a union and then from
   /// there we can have additional unnamed fields.
-  llvm::DIGlobalVariable *CollectAnonRecordDeclsForHeterogeneousDwarf(
+  llvm::DIGlobalVariable *CollectAnonRecordDeclsForHeterogeneousDwarfDIExpr(
+      const RecordDecl *RD, llvm::DIFile *Unit, unsigned LineNo,
+      StringRef LinkageName, llvm::dwarf::MemorySpace MS,
+      llvm::GlobalVariable *Var, llvm::DIScope *DContext);
+
+  /// Return a global variable that represents one of the collection of global
+  /// variables created for an anonmyous union (-gheterogeneous-dwarf).
+  ///
+  /// Recursively collect all of the member fields of a global
+  /// anonymous decl and create static variables for them. The first
+  /// time this is called it needs to be on a union and then from
+  /// there we can have additional unnamed fields.
+  llvm::DIGlobalVariableExpression *
+  CollectAnonRecordDeclsForHeterogeneousDwarfDIExpression(
       const RecordDecl *RD, llvm::DIFile *Unit, unsigned LineNo,
       StringRef LinkageName, llvm::dwarf::MemorySpace MS,
       llvm::GlobalVariable *Var, llvm::DIScope *DContext);

clang/lib/CodeGen/CGDebugInfo.h

@@ -985,7 +985,7 @@ void CodeGenModule::Release() {
     // We support a single version in the linked module. The LLVM
     // parser will drop debug info with a different version number
     // (and warn about it, too).
-    if (CodeGenOpts.HeterogeneousDwarf) {
+    if (CodeGenOpts.isHeterogeneousDwarfDIExpr()) {
       getModule().addModuleFlag(llvm::Module::Override, "Debug Info Version",
                                 llvm::DEBUG_METADATA_VERSION_HETEROGENEOUS_DWARF);
     } else {

clang/lib/CodeGen/CodeGenModule.cpp

@@ -4590,9 +4590,30 @@ renderDebugOptions(const ToolChain &TC, const Driver &D, const llvm::Triple &T,
   bool EmitDwarfForAMDGCN = EmitDwarf && T.isAMDGCN();
   if (EmitDwarfForAMDGCN)
     CmdArgs.append({"-mllvm", "-amdgpu-spill-cfi-saved-regs"});
-  if (Args.hasFlag(options::OPT_gheterogeneous_dwarf,
-                   options::OPT_gno_heterogeneous_dwarf, EmitDwarfForAMDGCN))
-    CmdArgs.push_back("-gheterogeneous-dwarf");
+  if (Arg *A = Args.getLastArg(options::OPT_gheterogeneous_dwarf_EQ)) {
+    A->render(Args, CmdArgs);
+  } else if (EmitDwarfForAMDGCN) {
+#ifndef NDEBUG
+    // There doesn't seem to be a straightforward way to "render" an option
+    // acquired from the OptTable into a string we can append to CmdArgs.
+    // All of the logic is buried in "accept" which works directly in terms
+    // of an ArgList.
+    //
+    // Instead, assert that the static string we are adding to CmdArgs has
+    // the same shape as what a bare -gheterogeneous-dwarf would alias to
+    // if the user has provided it in ArgList.
+    const Option GHeterogeneousDwarf =
+        getDriverOptTable().getOption(options::OPT_gheterogeneous_dwarf);
+    const Option Aliased = GHeterogeneousDwarf.getAlias();
+    assert(Aliased.isValid() && "gheterogeneous-dwarf must be an alias");
+    assert(Aliased.getName() == "gheterogeneous-dwarf=" &&
+           "gheterogeneous-dwarf must alias gheterogeneous-dwarf=");
+    assert(StringRef(GHeterogeneousDwarf.getAliasArgs()) == "diexpr" &&
+           GHeterogeneousDwarf.getAliasArgs()[strlen("diexpr") + 1] == '\0' &&
+           "gheterogeneous-dwarf must alias gheterogeneous-dwarf=diexpr");
+#endif
+    CmdArgs.push_back("-gheterogeneous-dwarf=diexpr");
+  }
 
   // This controls whether or not we perform JustMyCode instrumentation.
   if (Args.hasFlag(options::OPT_fjmc, options::OPT_fno_jmc, false)) {

clang/lib/Driver/ToolChains/Clang.cpp

@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -fblocks -debug-info-kind=limited -gheterogeneous-dwarf -emit-llvm -disable-llvm-verifier -o - %s | FileCheck %s
-// RUN: %clang_cc1 -DDEAD_CODE -fblocks -debug-info-kind=limited -gheterogeneous-dwarf -emit-llvm -disable-llvm-verifier -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fblocks -debug-info-kind=limited -gheterogeneous-dwarf=diexpr -emit-llvm -disable-llvm-verifier -o - %s | FileCheck %s
+// RUN: %clang_cc1 -DDEAD_CODE -fblocks -debug-info-kind=limited -gheterogeneous-dwarf=diexpr -emit-llvm -disable-llvm-verifier -o - %s | FileCheck %s
 
 typedef void (^BlockTy)();
 void escapeFunc(BlockTy);

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

@@ -1,12 +1,12 @@
-// RUN: %clang_cc1 -D ARG_TYPE=int -D PTR_ARG='&g' -D VAL_ARG=g -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf %s -o - | FileCheck --check-prefix=INT-ADDROF-VAL %s
-// RUN: %clang_cc1 -D ARG_TYPE=int -D PTR_ARG='&g' -D VAL_ARG=0 -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf %s -o - | FileCheck --check-prefix=INT-ADDROF-NOVAL %s
-// RUN: %clang_cc1 -D ARG_TYPE=int -D PTR_ARG=0 -D VAL_ARG=g -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf %s -o - | FileCheck --check-prefix=INT-NOADDROF-VAL %s
-// RUN: %clang_cc1 -D ARG_TYPE=int -D PTR_ARG=0 -D VAL_ARG=0 -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf %s -o - | FileCheck --check-prefix=INT-NOADDROF-NOVAL %s
+// RUN: %clang_cc1 -D ARG_TYPE=int -D PTR_ARG='&g' -D VAL_ARG=g -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf=diexpr %s -o - | FileCheck --check-prefix=INT-ADDROF-VAL %s
+// RUN: %clang_cc1 -D ARG_TYPE=int -D PTR_ARG='&g' -D VAL_ARG=0 -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf=diexpr %s -o - | FileCheck --check-prefix=INT-ADDROF-NOVAL %s
+// RUN: %clang_cc1 -D ARG_TYPE=int -D PTR_ARG=0 -D VAL_ARG=g -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf=diexpr %s -o - | FileCheck --check-prefix=INT-NOADDROF-VAL %s
+// RUN: %clang_cc1 -D ARG_TYPE=int -D PTR_ARG=0 -D VAL_ARG=0 -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf=diexpr %s -o - | FileCheck --check-prefix=INT-NOADDROF-NOVAL %s
 //
-// RUN: %clang_cc1 -D ARG_TYPE=float -D PTR_ARG='&g' -D VAL_ARG=g -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf %s -o - | FileCheck --check-prefix=FLOAT-ADDROF-VAL %s
-// RUN: %clang_cc1 -D ARG_TYPE=float -D PTR_ARG='&g' -D VAL_ARG=0 -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf %s -o - | FileCheck --check-prefix=FLOAT-ADDROF-NOVAL %s
-// RUN: %clang_cc1 -D ARG_TYPE=float -D PTR_ARG=0 -D VAL_ARG=g -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf %s -o - | FileCheck --check-prefix=FLOAT-NOADDROF-VAL %s
-// RUN: %clang_cc1 -D ARG_TYPE=float -D PTR_ARG=0 -D VAL_ARG=0 -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf %s -o - | FileCheck --check-prefix=FLOAT-NOADDROF-NOVAL %s
+// RUN: %clang_cc1 -D ARG_TYPE=float -D PTR_ARG='&g' -D VAL_ARG=g -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf=diexpr %s -o - | FileCheck --check-prefix=FLOAT-ADDROF-VAL %s
+// RUN: %clang_cc1 -D ARG_TYPE=float -D PTR_ARG='&g' -D VAL_ARG=0 -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf=diexpr %s -o - | FileCheck --check-prefix=FLOAT-ADDROF-NOVAL %s
+// RUN: %clang_cc1 -D ARG_TYPE=float -D PTR_ARG=0 -D VAL_ARG=g -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf=diexpr %s -o - | FileCheck --check-prefix=FLOAT-NOADDROF-VAL %s
+// RUN: %clang_cc1 -D ARG_TYPE=float -D PTR_ARG=0 -D VAL_ARG=0 -emit-llvm -debug-info-kind=standalone -gheterogeneous-dwarf=diexpr %s -o - | FileCheck --check-prefix=FLOAT-NOADDROF-NOVAL %s
 
 // INT-ADDROF-VAL: @g = internal constant i32 1, align 4{{$}}
 // INT-ADDROF-VAL-DAG: !llvm.dbg.retainedNodes = !{![[#LIFETIME:]]}

clang/test/CodeGen/debug-info-global-constant-heterogeneous-dwarf.c

@@ -1,4 +1,4 @@
-// RUN: %clang -g -gheterogeneous-dwarf -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang -g -gheterogeneous-dwarf=diexpr -S -emit-llvm -o - %s | FileCheck %s
 // RUN: %clang -S -emit-llvm -o - %s | FileCheck %s --check-prefix=NO_DEBUG
 int main (void) {
   return 0;

clang/test/CodeGen/debug-info-version-heterogeneous-dwarf.c

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple nvptx-unknown-unknown -emit-llvm -fcuda-is-device -debug-info-kind=limited -gheterogeneous-dwarf -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple nvptx-unknown-unknown -emit-llvm -fcuda-is-device -debug-info-kind=limited -gheterogeneous-dwarf=diexpr -o - %s | FileCheck %s
 // CHECK-DAG: !DIGlobalVariable(name: "GlobalShared", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true, memorySpace: DW_MSPACE_LLVM_group)
 // CHECK-DAG: !DIGlobalVariable(name: "GlobalDevice", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true, memorySpace: DW_MSPACE_LLVM_global)
 // CHECK-DAG: !DIGlobalVariable(name: "GlobalConstant", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: false, isDefinition: true, memorySpace: DW_MSPACE_LLVM_constant)

clang/test/CodeGenCUDA/debug-info-memory-space.cu

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x c++ -emit-llvm -fcuda-is-device -debug-info-kind=limited -gheterogeneous-dwarf -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x c++ -emit-llvm -fcuda-is-device -debug-info-kind=limited -gheterogeneous-dwarf=diexpr -o - %s | FileCheck %s
 
 struct S0 {
   unsigned int x : 16;

clang/test/CodeGenCXX/heterogeneous-debug-info-structured-binding-bitfield.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x c++ -emit-llvm -fcuda-is-device -debug-info-kind=limited -gheterogeneous-dwarf -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x c++ -emit-llvm -fcuda-is-device -debug-info-kind=limited -gheterogeneous-dwarf=diexpr -o - %s | FileCheck %s
 //
 // first def for 'a'
 // CHECK-LABEL: @_Z1fv()

clang/test/CodeGenCXX/heterogeneous-debug-info-structured-binding.cpp

@@ -1,5 +1,5 @@
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -debug-info-kind=limited -gheterogeneous-dwarf -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -debug-info-kind=limited -gheterogeneous-dwarf=diexpr -o - %s | FileCheck %s
 
 
 #define __device__ __attribute__((device))

clang/test/CodeGenHIP/debug-info-address-class-heterogeneous-dwarf.hip

@@ -1,5 +1,5 @@
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -O0 -debug-info-kind=limited -gheterogeneous-dwarf -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
+// RUN: %clang_cc1 -O0 -debug-info-kind=limited -gheterogeneous-dwarf=diexpr -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
 
 // Notes:
 // * There is no test involving transparent_union, as this isn't supported in

clang/test/CodeGenHIP/debug-info-amdgcn-abi-heterogeneous-dwarf.hip

@@ -1,5 +1,5 @@
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -debug-info-kind=limited -gheterogeneous-dwarf -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -debug-info-kind=limited -gheterogeneous-dwarf=diexpr -o - %s | FileCheck %s
 
 #define __device__ __attribute__((device))
 

clang/test/CodeGenHIP/debug-info-anonymous-union-heterogeneous-dwarf.hip

@@ -0,0 +1,12 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device -debug-info-kind=limited -gheterogeneous-dwarf -o - %s | FileCheck %s
+
+// Check that -gheterogeneous-dwarf without an `=OPTION` suffix remains valid
+// and aliases the new default. This is needed for transitioning flang-legacy
+// as it depends on the -cc1 interface.
+
+// CHECK: call void @llvm.dbg.def
+// CHECK: !DIExpr(
+__attribute__((device)) void kernel1(int Arg) {
+  int FuncVar;
+}