[XLA:GPU] Mark loads as invariant where appropriate.

If we read a value within an HLO that isn't modified by that HLO, mark
it as invariant in LLVM IR.

LLVM can perform more aggressive optimizations on invariant loads, but I
don't expect this will help much in our case, because XLA already emits
pretty aggressive noalias information on loads and stores.

However, on nvidia GPUs, marking loads as invariant has the additional
benefit of allowing LLVM to lower the load as ld.global.nc (equivalent
to CUDA's __ldg).  This instruction uses a special cache on the GPU, and
it's usually faster than a regular load.

PiperOrigin-RevId: 175774979