[NVPTX] remove unnecessary named metadata update that happens to break debug info.

[tra]

Summary

NVPTXGenericToNVVM pass updates named metadata by dropping all references to its arguments and uses addOperand() to add replacement entries. The problem is that there may bemetadata references to the argument we've attempted to replace and those would still refer to the old argument. When that happens we end up with metadata in inconsistent state which has both old and new metadata.

When following code is compiled with debug info, IR faill verification after NVPTXGenericToNVVM pass, because we end up with duplicate DICompileUnit metadata. Included test case demonstrates this when run with unpatched opt.

__device__ void bar(char* arg);
__global__ void foo() {
  bar("XYZ");
}

As far as I can tell, explicitly updating named metadata is not necessary -- updating metadata nodes for variables we've changed appears to maintain metadata consistency. I didn't manage to find a good way to verify it, though. If anyone has suggestions how to reference variable from named metadata, I can add it to the new test.

[dblaikie]

"As far as I can tell, explicitly updating named metadata is not necessary -- updating metadata nodes for variables we've changed appears to maintain metadata consistency. I didn't manage to find a good way to verify it, though. If anyone has suggestions how to reference variable from named metadata, I can add it to the new test."

Not sure I understand what you were and were not able to verify/test. Could you explain that more/differently?

The test case looks fairly reasonable to me - but I also wouldn't mind an explanation of all the features (why the artificial/named metadata node? I would expect to just match the CU/Globals/global metadata nodes, and check that the reference (as you have) uses whatever is necessary (the addrspace casts, etc) to still reference that global)

[tra]

In D22837#502775, @dblaikie wrote:

"As far as I can tell, explicitly updating named metadata is not necessary -- updating metadata nodes for variables we've changed appears to maintain metadata consistency. I didn't manage to find a good way to verify it, though. If anyone has suggestions how to reference variable from named metadata, I can add it to the new test."

Not sure I understand what you were and were not able to verify/test. Could you explain that more/differently?

I've verified that DI metadata is consistent after the pass. I.e. llvm.dbg.cu->DICompileUnit:globals->!3->DIGlobalVariable->correct reference to @static_var. Implicitly IR verification pass also verified that there's no leftover DICompileUnits hanging around which was the case before the patch.

What I'm not sure about is whether named metadata will always be handled correctly.
As far as I can tell, named metadata is pretty much read-only. You can only replace it with a new temporary instance.
Removed code was attempting to manually update named metadata if it happened to reference (directly or not) any IR that the pass changed.

Currently all named metadata is few references away from IR that gets changed. Regular metadata nodes can be modified (all?, some?), so changing a leaf node does not necessarily require replacement of the metadata node that references it. Updating single metadata node appears to be sufficient and does not require to propagate changes further.

However, if there's a way for a global node to refer to IR, that may need explicit replacement of such node.

The test case looks fairly reasonable to me - but I also wouldn't mind an explanation of all the features (why the artificial/named metadata node? I would expect to just match the CU/Globals/global metadata nodes, and check that the reference (as you have) uses whatever is necessary (the addrspace casts, etc) to still reference that global)

Artificial named node is a placeholder for a named metadata node that would have something that is modified by the pass and thus may need to be replaced. So far I've failed to create such metadata.

LLVM language reference says that named metadata can only have node operands, which guarantees at least one level of indirection, so named node can't refer to IR directly (i.e. it can't have DIGlobalVariable operands.). This simplifies things somehow.

What I need to know is -- when a regular metadata node is modified, am I guaranteed that:
a) either modification happend in-place, without replacing the node,
b) or, if the node is replaced, all metadata nodes referencing it (including named metadata nodes) are updated appropriately.

If that is true, then we're good. It appears to be the case, but I can't tell if I'm just missing a test case that would prove me wrong.

I guess I can try to see what happens when I apply ReplaceAllUsesWith() to a node used as an operand of named metadata. If that works correctly (i.e it would update named metadata operand to point to the new node), then I'm in the clear. If it does not, then I'll need to figure out how such case should be handled.

[tra]

Removed artificial named metadata node.
Added checks to verify DI consistency.

[dblaikie]

test/CodeGen/NVPTX/generic-to-nvvm-ir.ll

49 ↗

(On Diff #66509)

Is this not tested anywhere else?

57 ↗

(On Diff #66509)

This should probably be CHECK-SAME? & maybe just check the name of the DISubprogram for good measure? (same as you did with the global variable)

[tra]

Added few more metadata checks as dblaikie@ suggested.

[tra]

Removed mentions of named.md from comments.

llvm/trunk/lib/Target/NVPTX/NVPTXGenericToNVVM.cpp

Value *getOrInsertCVTA(Module *M, Function *F, GlobalVariable *GV,
                          IRBuilder<> &Builder);
   Value *remapConstant(Module *M, Function *F, Constant *C,
                        IRBuilder<> &Builder);
   Value *remapConstantVectorOrConstantAggregate(Module *M, Function *F,
                                                 Constant *C,
                                                 IRBuilder<> &Builder);
   Value *remapConstantExpr(Module *M, Function *F, ConstantExpr *C,
                            IRBuilder<> &Builder);
-  void remapNamedMDNode(ValueToValueMapTy &VM, NamedMDNode *N);
 
   typedef ValueMap<GlobalVariable *, GlobalVariable *> GVMapTy;
   typedef ValueMap<Constant *, Value *> ConstantToValueMapTy;
   GVMapTy GVMap;
   ConstantToValueMapTy ConstantToValueMap;
 };
 } // end namespace
 
for (Module::iterator I = M.begin(), E = M.end(); I != E; ++I) {
     ConstantToValueMap.clear();
   }
 
   // Copy GVMap over to a standard value map.
   ValueToValueMapTy VM;
   for (auto I = GVMap.begin(), E = GVMap.end(); I != E; ++I)
     VM[I->first] = I->second;
 
-  // Walk through the metadata section and update the debug information
-  // associated with the global variables in the default address space.
-  for (NamedMDNode &I : M.named_metadata()) {
-    remapNamedMDNode(VM, &I);
-  }
-
   // Walk through the global variable  initializers, and replace any use of
   // original global variables in GVMap with a use of the corresponding copies
   // in GVMap.  The copies need to be bitcast to the original global variable
   // types, as we cannot use cvta in global variable initializers.
   for (GVMapTy::iterator I = GVMap.begin(), E = GVMap.end(); I != E;) {
     GlobalVariable *GV = I->first;
     GlobalVariable *NewGV = I->second;
 
default:
     // UnaryConstantExpr
     if (Instruction::isCast(Opcode)) {
       return Builder.CreateCast(Instruction::CastOps(C->getOpcode()),
                                 NewOperands[0], C->getType());
     }
     llvm_unreachable("GenericToNVVM encountered an unsupported ConstantExpr");
   }
 }
-
-void GenericToNVVM::remapNamedMDNode(ValueToValueMapTy &VM, NamedMDNode *N) {
-
-  bool OperandChanged = false;
-  SmallVector<MDNode *, 16> NewOperands;
-  unsigned NumOperands = N->getNumOperands();
-
-  // Check if any operand is or contains a global variable in  GVMap, and thus
-  // converted to another value.
-  for (unsigned i = 0; i < NumOperands; ++i) {
-    MDNode *Operand = N->getOperand(i);
-    MDNode *NewOperand = MapMetadata(Operand, VM);
-    OperandChanged |= Operand != NewOperand;
-    NewOperands.push_back(NewOperand);
-  }
-
-  // If none of the operands has been modified, return immediately.
-  if (!OperandChanged) {
-    return;
-  }
-
-  // Replace the old operands with the new operands.
-  N->dropAllReferences();
-  for (SmallVectorImpl<MDNode *>::iterator I = NewOperands.begin(),
-                                           E = NewOperands.end();
-       I != E; ++I) {
-    N->addOperand(*I);
-  }
-}

llvm/trunk/test/CodeGen/NVPTX/generic-to-nvvm-ir.ll

+; Verify functionality of NVPTXGenericToNVVM.cpp pass.
+;
+; RUN: opt < %s -march nvptx64 -S -generic-to-nvvm -verify-debug-info | FileCheck %s
+
+target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
+target triple = "nvptx64-nvidia-cuda"
+
+; Generic space variables should be converted to global space AKA addrspace(1).
+; CHECK-DAG: @static_var = {{.*}}addrspace(1)
+@static_var = externally_initialized global i8 0, align 1
+; CHECK-DAG: @.str = {{.*}}addrspace(1)
+@.str = private unnamed_addr constant [4 x i8] c"XXX\00", align 1
+
+; Function Attrs: convergent
+define void @func() !dbg !8 {
+;CHECK-LABEL: @func()
+;CHECK-SAME: !dbg [[FUNCNODE:![0-9]+]]
+entry:
+; References to the variables must be converted back to generic address space via llvm intrinsic call
+; CHECK-DAG: call i8* @llvm.nvvm.ptr.global.to.gen.p0i8.p1i8({{.*}} addrspace(1)* @.str
+  %0 = load i8, i8* getelementptr inbounds ([4 x i8], [4 x i8]* @.str, i64 0, i64 0), align 1
+  call void @extfunc(i8 signext %0)
+; CHECK-DAG: call i8* @llvm.nvvm.ptr.global.to.gen.p0i8.p1i8(i8 addrspace(1)* @static_var
+  %1 = load i8, i8* @static_var, align 1
+  call void @extfunc(i8 signext %1)
+  ret void
+; CHECK: ret void
+}
+
+declare void @extfunc(i8 signext)
+
+!llvm.dbg.cu = !{!0}
+; CHECK: !llvm.dbg.cu = !{[[CUNODE:![0-9]+]]}
+!llvm.module.flags = !{!6, !7}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus, file: !1,
+      producer: "clang version 4.0.0",
+      isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, enums: !2, globals: !3)
+; CHECK: [[CUNODE]] = distinct !DICompileUnit({{.*}} globals: [[GLOBALSNODE:![0-9]+]]
+!1 = !DIFile(filename: "foo.cu", directory: "/usr/local/google/home/tra/work/llvm/build/gpu/debug")
+!2 = !{}
+!3 = !{!4}
+; Find list of global variables and make sure it's the one used by DICompileUnit
+; CHECK: [[GLOBALSNODE]] = !{[[GVNODE:![0-9]+]]}
+!4 = distinct !DIGlobalVariable(name: "static_var", scope: !0, file: !1, line: 2, type: !5, isLocal: false,
+               isDefinition: true, variable: i8* @static_var)
+; Debug info must also be updated to reflect new address space.
+; CHECK: [[GVNODE]] = distinct !DIGlobalVariable(name: "static_var"
+; CHECK-SAME: scope: [[CUNODE]]
+; CHECK-SAME: type: [[TYPENODE:![0-9]+]]
+; CHECK-SAME: variable: i8* addrspacecast (i8 addrspace(1)* @static_var to i8*)
+!5 = !DIBasicType(name: "char", size: 8, align: 8, encoding: DW_ATE_signed_char)
+; CHECK: [[TYPENODE]] = !DIBasicType(name: "char"
+!6 = !{i32 2, !"Dwarf Version", i32 4}
+!7 = !{i32 2, !"Debug Info Version", i32 3}
+!8 = distinct !DISubprogram(name: "foo", linkageName: "func",
+      scope: !1, file: !1, line: 3, type: !9, isLocal: false, isDefinition: true, scopeLine: 3,
+      flags: DIFlagPrototyped, isOptimized: false, unit: !0, variables: !2)
+; CHECK: [[FUNCNODE]] = distinct !DISubprogram(name: "foo",
+; CHECK-SAME: type: [[STYPENODE:![0-9]+]]
+; CHECK-SAME: unit: [[CUNODE]],
+!9 = !DISubroutineType(types: !10)
+; CHECK: [[STYPENODE]] = !DISubroutineType
+!10 = !{null}