`L1::no_allocate` is not safe to load volatile data

for the ptx instruction `ld.global.nc.L1::no_allocate.L2::256B` you mentioned, on devices which global memory is cached in L1 by default, such as Volta to Blackwell (sm70+), it's equivalent to `ld.global.L1::no_allocate.L2::256B`.

`L1::no_allocate` means the memory instruction doesn't allocate L1 cache line or check 'hit on miss' (the address is on the same cache line with a previous cache missed memory request), so access the data in a stale L1 cache line will not fetch the cache line from L2 or DRAM, and can not get the updated data from other SM.