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

[David Majnemer via llvm-commits]

Author: majnemer
Date: Wed Dec 21 18:51:59 2016
New Revision: 290294

URL: http://llvm.org/viewvc/llvm-project?rev=290294&view=rev
Log:
[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.

Modified:
    llvm/trunk/lib/Target/NVPTX/NVVMIntrRange.cpp
    llvm/trunk/test/CodeGen/NVPTX/intrinsic-old.ll

Modified: llvm/trunk/lib/Target/NVPTX/NVVMIntrRange.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVVMIntrRange.cpp?rev=290294&r1=290293&r2=290294&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVVMIntrRange.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVVMIntrRange.cpp Wed Dec 21 18:51:59 2016
@@ -65,6 +65,10 @@ INITIALIZE_PASS(NVVMIntrRange, "nvvm-int
 // 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[] = {

Modified: llvm/trunk/test/CodeGen/NVPTX/intrinsic-old.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/intrinsic-old.ll?rev=290294&r1=290293&r2=290294&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/intrinsic-old.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/intrinsic-old.ll Wed Dec 21 18:51:59 2016
@@ -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}