summaryrefslogtreecommitdiffstats
path: root/llvm/lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp
Commit message (Collapse)AuthorAgeFilesLines
* NVPTX: Remove implicit ilist iterator conversions, NFCDuncan P. N. Exon Smith2015-10-201-4/+3
| | | | llvm-svn: 250779
* Fix some comment typos.Benjamin Kramer2015-08-081-1/+1
| | | | llvm-svn: 244402
* [NVPTX] convert pointers in byval kernel arguments to globalJingyue Wu2015-07-311-20/+81
| | | | | | | | | | | | | | | | | | | | | | | | | 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-0/+4
| | | | | | | | | | | | | | | | 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/+170
| | | | | | | | | 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-170/+0
| | | | | | llc crashed for NVPTX backend llvm-svn: 239094
* [NVPTX] kernel pointer arguments point to the global address spaceJingyue Wu2015-06-041-0/+170
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