summaryrefslogtreecommitdiffstats
path: root/llvm/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll
Commit message (Collapse)AuthorAgeFilesLines
* [NVPTX] Improve lowering of byval args of device functions.Artem Belevich2016-07-201-7/+25
| | | | | | | | | | | | Avoid unnecessary spills of byval arguments of device functions to local space on SASS level and subsequent pointer conversion to generic address space that follows. Instead, make a local copy in IR, provide a way to access arguments directly, and let LLVM optimize the copy away when possible. Differential Review: https://reviews.llvm.org/D21421 llvm-svn: 276153
* Revert r273313 "[NVPTX] Improve lowering of byval args of device functions."Artem Belevich2016-06-291-25/+7
| | | | | | The change causes llvm crash in some unoptimized builds. llvm-svn: 274163
* [NVPTX] Improve lowering of byval args of device functions.Artem Belevich2016-06-211-7/+25
| | | | | | | | | | | | | | | | | | | | | Avoid unnecessary spills of such vars to local space on SASS level and pointer space conversion. Instead, make a local copy with appropriate addrspacecasts and let LLVM optimize them away when possible. This allows loading value of the argument using [symbol+offset] instead of converting argument to general space pointer and using it for indexing (which also implicitly converts param space pointer to local space one on SASS level and triggers copying of argument into local space in the process). This reduces call overhead, uses less registers and reduces overall SASS size by 2-4%. Differential Review: http://reviews.llvm.org/D21421 llvm-svn: 273313
* [NVPTX] convert pointers in byval kernel arguments to globalJingyue Wu2015-07-311-2/+18
| | | | | | | | | | | | | | | | | | | | | | | | | Summary: For example, in struct S { int *x; int *y; }; __global__ void foo(S s) { int *b = s.y; // use b } "b" is guaranteed to point to global. NVPTX should emit ld.global/st.global for accessing "b". Reviewers: jholewinski Subscribers: llvm-commits, jholewinski Differential Revision: http://reviews.llvm.org/D11505 llvm-svn: 243790
* [NVPTX] noop when kernel pointers are already globalJingyue Wu2015-06-261-1/+12
| | | | | | | | | | | | | | | | 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
* [NVPTX] roll forward r239082Jingyue Wu2015-06-041-0/+20
| | | | | | | | | NVPTXISelDAGToDAG translates "addrspacecast to param" to NVPTX::nvvm_ptr_gen_to_param Added an llc test in bug21465. llvm-svn: 239100
* Revert r239082Jingyue Wu2015-06-041-20/+0
| | | | | | llc crashed for NVPTX backend llvm-svn: 239094
* [NVPTX] kernel pointer arguments point to the global address spaceJingyue Wu2015-06-041-0/+20
Summary: With this patch, NVPTXLowerKernelArgs converts a kernel pointer argument to a pointer in the global address space. This change, along with NVPTXFavorNonGenericAddrSpaces, allows the NVPTX backend to emit ld.global.* and st.global.* for accessing kernel pointer arguments. Minor changes: 1. refactor: extract function convertToPointerInAddrSpace 2. fix a bug in the test case in bug21465.ll Test Plan: lower-kernel-ptr-arg.ll Reviewers: eliben, meheff, jholewinski Reviewed By: jholewinski Subscribers: wengxt, jholewinski, llvm-commits Differential Revision: http://reviews.llvm.org/D10154 llvm-svn: 239082
OpenPOWER on IntegriCloud