4 Commits

Author SHA1 Message Date
David Majnemer
f35020be62 [NVVMIntrRange] Only set range metadata if none is already present
The range metadata inserted by NVVMIntrRange is pessimistic, range
metadata already present could be more precise.

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@290294 91177308-0d34-0410-b5e6-96231b3b80d8
2016-12-22 00:51:59 +00:00
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