llvm.org GIT mirror llvm / a70d990
[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 Jingyue Wu 4 years ago
2 changed file(s) with 16 addition(s) and 1 deletion(s). Raw diff Collapse all Expand all
131131 assert(!Arg->hasByValAttr() &&
132132 "byval params should be handled by handleByValParam");
133133
134 // Do nothing if the argument already points to the global address space.
135 if (Arg->getType()->getPointerAddressSpace() == ADDRESS_SPACE_GLOBAL)
136 return;
137
134138 Instruction *FirstInst = Arg->getParent()->getEntryBlock().begin();
135139 Instruction *ArgInGlobal = new AddrSpaceCastInst(
136140 Arg, PointerType::get(Arg->getType()->getPointerElementType(),
1515 ret void
1616 }
1717
18 !nvvm.annotations = !{!0}
18 define void @kernel2(float addrspace(1)* %input, float addrspace(1)* %output) {
19 ; CHECK-LABEL: .visible .entry kernel2(
20 ; CHECK-NOT: cvta.to.global.u64
21 %1 = load float, float addrspace(1)* %input, align 4
22 ; CHECK: ld.global.f32
23 store float %1, float addrspace(1)* %output, align 4
24 ; CHECK: st.global.f32
25 ret void
26 }
27
28 !nvvm.annotations = !{!0, !1}
1929 !0 = !{void (float*, float*)* @kernel, !"kernel", i32 1}
30 !1 = !{void (float addrspace(1)*, float addrspace(1)*)* @kernel2, !"kernel", i32 1}