From 3203818bf7df5ff7b0b01b036cfeeb090ab9e596 Mon Sep 17 00:00:00 2001 From: Jingyue Wu Date: Fri, 26 Jun 2015 22:35:43 +0000 Subject: [PATCH] [NVPTX] noop when kernel pointers are already global Summary: Some front ends make kernel pointers global already. In that case, handlePointerParams does nothing. Test Plan: more tests in lower-kernel-ptr-arg.ll Reviewers: grosser Subscribers: jholewinski, llvm-commits Differential Revision: http://reviews.llvm.org/D10779 llvm-svn: 240849 --- llvm/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp | 4 ++++ llvm/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll | 13 ++++++++++++- 2 files changed, 16 insertions(+), 1 deletion(-) diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp b/llvm/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp index 24dcb12..b533f31 100644 --- a/llvm/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp @@ -132,6 +132,10 @@ void NVPTXLowerKernelArgs::handlePointerParam(Argument *Arg) { assert(!Arg->hasByValAttr() && "byval params should be handled by handleByValParam"); + // Do nothing if the argument already points to the global address space. + if (Arg->getType()->getPointerAddressSpace() == ADDRESS_SPACE_GLOBAL) + return; + Instruction *FirstInst = Arg->getParent()->getEntryBlock().begin(); Instruction *ArgInGlobal = new AddrSpaceCastInst( Arg, PointerType::get(Arg->getType()->getPointerElementType(), diff --git a/llvm/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll b/llvm/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll index 53220bd..0de72c4 100644 --- a/llvm/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll +++ b/llvm/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll @@ -16,5 +16,16 @@ define void @kernel(float* %input, float* %output) { ret void } -!nvvm.annotations = !{!0} +define void @kernel2(float addrspace(1)* %input, float addrspace(1)* %output) { +; CHECK-LABEL: .visible .entry kernel2( +; CHECK-NOT: cvta.to.global.u64 + %1 = load float, float addrspace(1)* %input, align 4 +; CHECK: ld.global.f32 + store float %1, float addrspace(1)* %output, align 4 +; CHECK: st.global.f32 + ret void +} + +!nvvm.annotations = !{!0, !1} !0 = !{void (float*, float*)* @kernel, !"kernel", i32 1} +!1 = !{void (float addrspace(1)*, float addrspace(1)*)* @kernel2, !"kernel", i32 1} -- 2.7.4