Commit Graph

3 Commits

Author SHA1 Message Date
Justin Bogner cde945cbcd NVPTX: Remove the legacy ptx intrinsics
- Rename the ptx.read.* intrinsics to nvvm.read.ptx.sreg.* - some but
  not all of these registers were already accessible via the nvvm
  name.
- Rename ptx.bar.sync nvvm.bar.sync, to match nvvm.bar0.

There's a fair amount of code motion here, but it's all very
mechanical.

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@274769 91177308-0d34-0410-b5e6-96231b3b80d8
2016-07-07 16:40:17 +00:00
Artem Belevich c1645b33e4 Init member structs in constructor.
Fixes build error on windows where MSVC does not
support list initialization inside member initializer list.

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@270877 91177308-0d34-0410-b5e6-96231b3b80d8
2016-05-26 17:29:20 +00:00
Artem Belevich bd2219ad6e [NVPTX] Added NVVMIntrRange pass
NVVMIntrRange adds !range metadata to calls of NVVM intrinsics
that return values within known limited range.

This allows LLVM to generate optimal code for indexing arrays
based on tid/ctaid which is a frequently used pattern in CUDA code.

Differential Revision: http://reviews.llvm.org/D20644

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@270872 91177308-0d34-0410-b5e6-96231b3b80d8
2016-05-26 17:02:56 +00:00