llvm-project

ce465594 - [NVPTX] Fix aggregate load/store lowering for (potentially) overlapping copies (#201177)

[NVPTX] Fix aggregate load/store lowering for (potentially) overlapping copies (#201177) NVPTXLowerAggrCopies lowers load/store pairs of large values into a loop of smaller copies. However, it was incorrectly assuming that the load/store pairs it found never alias. This patch adds an alias check. If the pointers may alias, we emit a memmov, which handles overlap correctly. CUDA reproducer: typedef char vec __attribute__((vector_size(256))); __global__ void boom(char *p) { *(vec *)(p + 8) = *(vec *)p; }