[PATCH] D11314: [NVPTX] make load on global readonly memory to use ldg

[Xuetian Weng]

wengxt created this revision.
wengxt added reviewers: jholewinski, jingyue.
wengxt added a subscriber: llvm-commits.
Herald added a subscriber: jholewinski.

As describe in [1], ld.global.nc may be used to load memory by nvcc when
__restrict__ is used and compiler can detect whether read-only data cache
is safe to use.

This patch will try to check whether ldg is safe to use and use them to
replace ld.global when possible.

Test Plan: test/CodeGen/NVPTX/load-with-non-coherent-cache.ll

http://reviews.llvm.org/D11314

Files:
  lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
  test/CodeGen/NVPTX/load-with-non-coherent-cache.ll

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D11314.30042.patch
Type: text/x-patch
Size: 11622 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20150717/decd42ba/attachment.bin>