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

Artem Belevich via llvm-commits llvm-commits at lists.llvm.org
Tue Jul 26 16:54:07 PDT 2016


tra created this revision.
tra added reviewers: dblaikie, jholewinski.
tra added a subscriber: llvm-commits.
Herald added a subscriber: jholewinski.

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.


https://reviews.llvm.org/D22837

Files:
  lib/Target/NVPTX/NVPTXGenericToNVVM.cpp
  test/CodeGen/NVPTX/generic-to-nvvm-ir.ll

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D22837.65629.patch
Type: text/x-patch
Size: 5247 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160726/172a295b/attachment.bin>


More information about the llvm-commits mailing list