[all-commits] [llvm/llvm-project] ce4655: [NVPTX] Fix aggregate load/store lowering for (pot...

[Justin Lebar via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: ce465594e2394ebe6894c6f0dfc9d0817b63c9da
      https://github.com/llvm/llvm-project/commit/ce465594e2394ebe6894c6f0dfc9d0817b63c9da
  Author: Justin Lebar <justin.lebar at gmail.com>
  Date:   2026-06-02 (Tue, 02 Jun 2026)

  Changed paths:
    M llvm/lib/Target/NVPTX/NVPTXLowerAggrCopies.cpp
    M llvm/test/CodeGen/NVPTX/lower-aggr-copies.ll

  Log Message:
  -----------
  [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;
    }



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications