[PATCH] D112466: [NVPTX] Drop memory references of LDG/LDU

[Artem Belevich via Phabricator via llvm-commits]

tra added a comment.

While `ldu` does indeed specify that it loads from read-only memory, I do not think we can treat `ld.global.nc` the same way.
PTX spec says `Load register variable d from the location specified by the source address operand a in the global state space, and optionally cache in non-coherent texture cache. Since the cache is non-coherent, the data should be read-only within the kernel's process.`

The way I read it -- it's a regular load that bypasses cache.  Unlike `ldu`, it does not specify that it's a read-only data. While read-only data will make it work correctly, it's not the only valid use case.
E.g. data may be changed from the host and two subsequent `ld.global.nc` will return different values.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D112466/new/

https://reviews.llvm.org/D112466