[PATCH] D21421: [NVPTX] Improve lowering of byval args of device functions.

[Artem Belevich via llvm-commits]

tra updated this revision to Diff 64422.
tra added a comment.
This revision is now accepted and ready to land.

ptxas bug workaround has been committed separately, so now it's purely an optimization.

Now byval arguments are handled the same way for both __global__ and __device__ functions:

- NVPTXLowerKernelArgs pass creates local copy which is addrspacecast'ed to param space.
- LowerFormalArguemts lowers them to MoveParams(param_symbol)
- SelectDirectAddr() recognizes addrspacecast([0->101] MoveParam(arg_symbol)) created above and returns arg_symbol.


https://reviews.llvm.org/D21421

Files:
  lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
  lib/Target/NVPTX/NVPTXISelLowering.cpp
  lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp
  test/CodeGen/NVPTX/bug21465.ll
  test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D21421.64422.patch
Type: text/x-patch
Size: 7129 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160718/ec4f0ffa/attachment.bin>