From 5fa7d48bb804dd507e64086636f2f725d4e3af71 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. llvm-svn: 290294 --- llvm/lib/Target/NVPTX/NVVMIntrRange.cpp | 4 ++++ llvm/test/CodeGen/NVPTX/intrinsic-old.ll | 10 ++++++++++ 2 files changed, 14 insertions(+) diff --git a/llvm/lib/Target/NVPTX/NVVMIntrRange.cpp b/llvm/lib/Target/NVPTX/NVVMIntrRange.cpp index b9c02c4..9c71a2e 100644 --- a/llvm/lib/Target/NVPTX/NVVMIntrRange.cpp +++ b/llvm/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/llvm/test/CodeGen/NVPTX/intrinsic-old.ll b/llvm/test/CodeGen/NVPTX/intrinsic-old.ll index daf83a8..4953af6 100644 --- a/llvm/test/CodeGen/NVPTX/intrinsic-old.ll +++ b/llvm/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.7.4