3 Commits

Author	SHA1	Message	Date
Artem Belevich	a28e598ebb	[NVPTX] Fixed vectorized LDG for f16. ... v2f16 is a special case in NVPTX. v4f16 may be loaded as a pair of v2f16 and that was not previously handled correctly by tryLDGLDU() Differential Revision: https://reviews.llvm.org/D45339 llvm-svn: 329456	2018-04-06 21:10:24 +00:00
Justin Lebar	e0550591a5	[NVPTX] Add tests that invariant vector loads get lowered to ld.global.nc. ... llvm-svn: 294082	2017-02-04 01:54:56 +00:00
Justin Lebar	6d6b11a4a6	[NVPTX] Use ldg for explicitly invariant loads. ... Summary: With this change (plus some changes to prevent !invariant from being clobbered within llvm), clang will be able to model the __ldg CUDA builtin as an invariant load, rather than as a target-specific llvm intrinsic. This will let the optimizer play with these loads -- specifically, we should be able to vectorize them in the load-store vectorizer. Reviewers: tra Subscribers: jholewinski, hfinkel, llvm-commits, chandlerc Differential Revision: https://reviews.llvm.org/D23477 llvm-svn: 281152	2016-09-11 01:39:04 +00:00