diff options
| -rw-r--r-- | llvm/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp | 4 | ||||
| -rw-r--r-- | llvm/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll | 13 |
2 files changed, 16 insertions, 1 deletions
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp b/llvm/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp index 24dcb122..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..0de72c4a 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} |