From f35020be62e1c819689a1622f360452e7374a42c Mon Sep 17 00:00:00 2001 From: David Majnemer Date: Thu, 22 Dec 2016 00:51:59 +0000 Subject: [PATCH] [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 --- lib/Target/NVPTX/NVVMIntrRange.cpp | 4 ++++ test/CodeGen/NVPTX/intrinsic-old.ll | 10 ++++++++++ 2 files changed, 14 insertions(+) diff --git a/lib/Target/NVPTX/NVVMIntrRange.cpp b/lib/Target/NVPTX/NVVMIntrRange.cpp index b9c02c43114..9c71a2ee165 100644 --- 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 daf83a87007..4953af6a141 100644 --- 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} -- 2.49.0