[NVVMIntrRange] Only set range metadata if none is already present

The range metadata inserted by NVVMIntrRange is pessimistic, range
metadata already present could be more precise.

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@290294 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/Target/NVPTX/NVVMIntrRange.cpp b/lib/Target/NVPTX/NVVMIntrRange.cpp
index b9c02c43114124cfc7680866b970c5660bddc511..9c71a2ee165b3c5c5689cbb57c42fdb49f475c4a 100644 (file)
--- a/lib/Target/NVPTX/NVVMIntrRange.cpp
+++ b/lib/Target/NVPTX/NVVMIntrRange.cpp
@@ -65,6 +65,10 @@ INITIALIZE_PASS(NVVMIntrRange, "nvvm-intr-range",
 // Adds the passed-in [Low,High) range information as metadata to the
 // passed-in call instruction.
 static bool addRangeMetadata(uint64_t Low, uint64_t High, CallInst *C) {
+  // This call already has range metadata, nothing to do.
+  if (C->getMetadata(LLVMContext::MD_range))
+    return false;
+
   LLVMContext &Context = C->getParent()->getContext();
   IntegerType *Int32Ty = Type::getInt32Ty(Context);
   Metadata *LowAndHigh[] = {

diff --git a/test/CodeGen/NVPTX/intrinsic-old.ll b/test/CodeGen/NVPTX/intrinsic-old.ll
index daf83a870075d46f5234b9584b49342f709d902a..4953af6a1413aeda3ef41f2577fb8cffd7fa409f 100644 (file)
--- a/test/CodeGen/NVPTX/intrinsic-old.ll
+++ b/test/CodeGen/NVPTX/intrinsic-old.ll
@@ -155,6 +155,13 @@ define ptx_device i32 @test_nctaid_x() {
        ret i32 %x
 }
 
+define ptx_device i32 @test_already_has_range_md() {
+; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.x;
+; RANGE: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x(), !range ![[ALREADY:[0-9]+]]
+       %x = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x(), !range !0
+       ret i32 %x
+}
+
 
 define ptx_device i32 @test_nctaid_w() {
 ; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.w;
@@ -311,6 +318,9 @@ declare i32 @llvm.nvvm.read.ptx.sreg.pm3()
 
 declare void @llvm.nvvm.bar.sync(i32 %i)
 
+!0 = !{i32 0, i32 19}
+; RANGE-DAG: ![[ALREADY]] = !{i32 0, i32 19}
+; RANGE-DAG: ![[BLK_IDX_XY]] = !{i32 0, i32 1024}
 ; RANGE-DAG: ![[BLK_IDX_XY]] = !{i32 0, i32 1024}
 ; RANGE-DAG: ![[BLK_IDX_Z]] = !{i32 0, i32 64}
 ; RANGE-DAG: ![[BLK_SIZE_XY]] = !{i32 1, i32 1025}