From a70d990f47315e9197a0d704884e980e3d3a2fa2 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 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@240849 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp | 4 ++++ test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll | 13 ++++++++++++- 2 files changed, 16 insertions(+), 1 deletion(-) diff --git a/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp b/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp index 24dcb122b94..b533f316d8a 100644 --- a/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp +++ b/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/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll b/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll index 53220bd905b..0de72c4a1ae 100644 --- a/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll +++ b/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.34.1