[llvm] [NVPTX] Fix aggregate load/store lowering for (potentially) overlapping copies (PR #201177)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Jun 2 11:56:48 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Justin Lebar (jlebar)
<details>
<summary>Changes</summary>
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;
}
---
Full diff: https://github.com/llvm/llvm-project/pull/201177.diff
2 Files Affected:
- (modified) llvm/lib/Target/NVPTX/NVPTXLowerAggrCopies.cpp (+50-11)
- (modified) llvm/test/CodeGen/NVPTX/lower-aggr-copies.ll (+118-1)
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAggrCopies.cpp b/llvm/lib/Target/NVPTX/NVPTXLowerAggrCopies.cpp
index afee989b5c466..3966085d27581 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAggrCopies.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAggrCopies.cpp
@@ -14,16 +14,19 @@
#include "NVPTXLowerAggrCopies.h"
#include "NVPTX.h"
+#include "llvm/Analysis/AliasAnalysis.h"
#include "llvm/Analysis/TargetTransformInfo.h"
#include "llvm/CodeGen/StackProtector.h"
#include "llvm/IR/Constants.h"
#include "llvm/IR/DataLayout.h"
#include "llvm/IR/Function.h"
+#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/IntrinsicInst.h"
#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/LLVMContext.h"
#include "llvm/IR/Module.h"
+#include "llvm/InitializePasses.h"
#include "llvm/Transforms/Utils/BasicBlockUtils.h"
#include "llvm/Transforms/Utils/LowerMemIntrinsics.h"
@@ -42,6 +45,7 @@ struct NVPTXLowerAggrCopies : public FunctionPass {
void getAnalysisUsage(AnalysisUsage &AU) const override {
AU.addPreserved<StackProtector>();
AU.addRequired<TargetTransformInfoWrapperPass>();
+ AU.addRequired<AAResultsWrapperPass>();
}
bool runOnFunction(Function &F) override;
@@ -63,6 +67,7 @@ bool NVPTXLowerAggrCopies::runOnFunction(Function &F) {
LLVMContext &Context = F.getParent()->getContext();
const TargetTransformInfo &TTI =
getAnalysis<TargetTransformInfoWrapperPass>().getTTI(F);
+ AAResults &AA = getAnalysis<AAResultsWrapperPass>().getAAResults();
// Collect all aggregate loads and mem* calls.
for (BasicBlock &BB : F) {
@@ -108,14 +113,43 @@ bool NVPTXLowerAggrCopies::runOnFunction(Function &F) {
ConstantInt *CopyLen =
ConstantInt::get(Type::getInt32Ty(Context), NumLoads);
- createMemCpyLoopKnownSize(/* ConvertedInst */ SI,
- /* SrcAddr */ SrcAddr, /* DstAddr */ DstAddr,
- /* CopyLen */ CopyLen,
- /* SrcAlign */ LI->getAlign(),
- /* DestAlign */ SI->getAlign(),
- /* SrcIsVolatile */ LI->isVolatile(),
- /* DstIsVolatile */ SI->isVolatile(),
- /* CanOverlap */ true, TTI);
+ LocationSize Size = LocationSize::precise(NumLoads);
+ if (AA.isNoAlias(MemoryLocation(SrcAddr, Size),
+ MemoryLocation(DstAddr, Size))) {
+ // No overlap: emit a plain memcpy loop. Expand the loop here (rather
+ // than emitting a memcpy intrinsic and letting the code below expand it)
+ // so we can pass CanOverlap = false; expandMemCpyAsLoop would
+ // conservatively assume overlap.
+ createMemCpyLoopKnownSize(/* ConvertedInst */ SI,
+ /* SrcAddr */ SrcAddr, /* DstAddr */ DstAddr,
+ /* CopyLen */ CopyLen,
+ /* SrcAlign */ LI->getAlign(),
+ /* DestAlign */ SI->getAlign(),
+ /* SrcIsVolatile */ LI->isVolatile(),
+ /* DstIsVolatile */ SI->isVolatile(),
+ /* CanOverlap */ false, TTI);
+ } else {
+ // May alias: lower as a memmove, which picks the copy direction at
+ // runtime. Emit the intrinsic here and let the loop below expand it.
+ //
+ // The pointers may alias even if they're in different address spaces
+ // (e.g. the generic addrspace may alias global). If they're in
+ // different addrspaces, cast to the generic space first, because
+ // expandMemMoveAsLoop needs to compare the pointer values to determine
+ // the copy direction.
+ IRBuilder<> Builder(SI);
+ unsigned SrcAS = LI->getPointerAddressSpace();
+ unsigned DstAS = SI->getPointerAddressSpace();
+ if (SrcAS != DstAS) {
+ PointerType *GenericPtrTy =
+ PointerType::get(Context, NVPTXAS::ADDRESS_SPACE_GENERIC);
+ SrcAddr = Builder.CreateAddrSpaceCast(SrcAddr, GenericPtrTy);
+ DstAddr = Builder.CreateAddrSpaceCast(DstAddr, GenericPtrTy);
+ }
+ MemCalls.push_back(cast<MemMoveInst>(Builder.CreateMemMove(
+ DstAddr, SI->getAlign(), SrcAddr, LI->getAlign(), CopyLen,
+ LI->isVolatile() || SI->isVolatile())));
+ }
SI->eraseFromParent();
LI->eraseFromParent();
@@ -140,9 +174,14 @@ bool NVPTXLowerAggrCopies::runOnFunction(Function &F) {
} // namespace
-INITIALIZE_PASS(NVPTXLowerAggrCopies, "nvptx-lower-aggr-copies",
- "Lower aggregate copies, and llvm.mem* intrinsics into loops",
- false, false)
+INITIALIZE_PASS_BEGIN(NVPTXLowerAggrCopies, "nvptx-lower-aggr-copies",
+ "Lower aggregate copies, and llvm.mem* intrinsics into loops",
+ false, false)
+INITIALIZE_PASS_DEPENDENCY(AAResultsWrapperPass)
+INITIALIZE_PASS_DEPENDENCY(TargetTransformInfoWrapperPass)
+INITIALIZE_PASS_END(NVPTXLowerAggrCopies, "nvptx-lower-aggr-copies",
+ "Lower aggregate copies, and llvm.mem* intrinsics into loops",
+ false, false)
FunctionPass *llvm::createLowerAggrCopies() {
return new NVPTXLowerAggrCopies();
diff --git a/llvm/test/CodeGen/NVPTX/lower-aggr-copies.ll b/llvm/test/CodeGen/NVPTX/lower-aggr-copies.ll
index 6c96eab1439fb..96951e7b2b04f 100644
--- a/llvm/test/CodeGen/NVPTX/lower-aggr-copies.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-aggr-copies.ll
@@ -1,5 +1,5 @@
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -O0 | FileCheck %s --check-prefix PTX
-; RUN: opt < %s -S -nvptx-lower-aggr-copies | FileCheck %s --check-prefix IR
+; RUN: opt < %s -S -nvptx-aa -nvptx-aa-wrapper -nvptx-lower-aggr-copies | FileCheck %s --check-prefix IR
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 -O0 | %ptxas-verify %}
; Verify that the NVPTXLowerAggrCopies pass works as expected - calls to
@@ -179,3 +179,120 @@ entry:
; PTX-NEXT: st.param.b64 [func_retval0
; PTX-NEXT: ret
}
+
+define void @aggr_loadstore_overlap_forward_copy(ptr %p) {
+entry:
+ %dst = getelementptr inbounds i8, ptr %p, i64 8
+ %v = load [128 x i8], ptr %p, align 1
+ store [128 x i8] %v, ptr %dst, align 1
+ ret void
+
+; A large aggregate load;store pair may have overlapping src/dst and is fully
+; defined (whole value read before any byte stored). It must be lowered to
+; overlap-safe memmove-style code (runtime src<dst direction check + backward
+; loop), not an unconditional forward copy loop.
+; IR-LABEL: @aggr_loadstore_overlap_forward_copy
+; IR: [[CMP:%[0-9a-zA-Z_]+]] = icmp ult ptr %p, %dst
+; IR: br i1 [[CMP]], label %memmove_bwd_loop, label %memmove_fwd_loop
+; IR: memmove_bwd_loop:
+; IR: %bwd_index = sub i32 {{%[0-9a-zA-Z_]+}}, 1
+; IR: memmove_fwd_loop:
+; IR: {{%[0-9a-zA-Z_]+}} = add i32 %fwd_index, 1
+
+; PTX-LABEL: .visible .func aggr_loadstore_overlap_forward_copy(
+; PTX: setp.ge.u64 %p{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}}
+; PTX: // %memmove_bwd_loop
+; PTX: // %memmove_fwd_loop
+}
+
+define void @aggr_loadstore_generic_global(ptr %g, ptr addrspace(1) %glob) {
+ %v = load [128 x i8], ptr %g, align 1
+ store [128 x i8] %v, ptr addrspace(1) %glob, align 1
+ ret void
+
+; The generic address space aliases every space, so a generic and a global
+; pointer may overlap and the copy must be direction-safe. The two pointers
+; live in different spaces, so both are cast to generic to make the runtime
+; comparison well defined before emitting the memmove-style loop.
+; IR-LABEL: @aggr_loadstore_generic_global
+; IR: [[GG:%[0-9a-zA-Z_]+]] = addrspacecast ptr addrspace(1) %glob to ptr
+; IR: [[CMP:%[0-9a-zA-Z_]+]] = icmp ult ptr %g, [[GG]]
+; IR: br i1 [[CMP]], label %memmove_bwd_loop, label %memmove_fwd_loop
+; IR: memmove_fwd_loop:
+
+; PTX-LABEL: .visible .func aggr_loadstore_generic_global(
+; PTX: cvta.global.u64
+; PTX: // %memmove_bwd_loop
+; PTX: // %memmove_fwd_loop
+}
+
+define void @aggr_loadstore_global_shared(ptr addrspace(1) %glob, ptr addrspace(3) %sh) {
+ %v = load [128 x i8], ptr addrspace(1) %glob, align 1
+ store [128 x i8] %v, ptr addrspace(3) %sh, align 1
+ ret void
+
+; Distinct non-generic address spaces (global vs shared) cannot overlap; this
+; fact comes from NVPTXAAResult (the RUN line adds nvptx-aa). So it lowers to a
+; plain forward copy loop with no runtime direction check, and the loads/stores
+; carry alias-scope/noalias metadata.
+; IR-LABEL: @aggr_loadstore_global_shared
+; IR-NOT: memmove_bwd_loop
+; IR-NOT: addrspacecast
+; IR: load i8, ptr addrspace(1) {{.*}}, !alias.scope
+; IR: store i8 {{.*}}, ptr addrspace(3) {{.*}}, !noalias
+
+; PTX-LABEL: .visible .func aggr_loadstore_global_shared(
+; PTX: // %static-memcpy
+; PTX-NOT: // %memmove_bwd_loop
+}
+
+define void @aggr_loadstore_shared_cluster(ptr addrspace(3) %sh, ptr addrspace(7) %clus) {
+ %v = load [128 x i8], ptr addrspace(3) %sh, align 1
+ store [128 x i8] %v, ptr addrspace(7) %clus, align 1
+ ret void
+
+; Distributed shared (addrspace 7) aliases shared (addrspace 3), so this pair
+; may overlap. The spaces differ, so both pointers are cast to generic for the
+; runtime comparison; cvta.shared / cvta.shared::cluster make that legal.
+; IR-LABEL: @aggr_loadstore_shared_cluster
+; IR: [[S:%[0-9a-zA-Z_]+]] = addrspacecast ptr addrspace(3) %sh to ptr
+; IR: [[C:%[0-9a-zA-Z_]+]] = addrspacecast ptr addrspace(7) %clus to ptr
+; IR: [[CMP:%[0-9a-zA-Z_]+]] = icmp ult ptr [[S]], [[C]]
+; IR: br i1 [[CMP]], label %memmove_bwd_loop, label %memmove_fwd_loop
+
+; PTX-LABEL: .visible .func aggr_loadstore_shared_cluster(
+; PTX: cvta.shared.u64
+; PTX: cvta.shared::cluster.u64
+; PTX: // %memmove_bwd_loop
+; PTX: // %memmove_fwd_loop
+}
+
+define void @aggr_loadstore_same_global(ptr addrspace(1) %a, ptr addrspace(1) %b) {
+ %v = load [128 x i8], ptr addrspace(1) %a, align 1
+ store [128 x i8] %v, ptr addrspace(1) %b, align 1
+ ret void
+
+; Same (non-generic) address space: may overlap, but the pointers are already
+; comparable, so no addrspacecast is introduced and the overlap-safe loop runs
+; directly in that space.
+; IR-LABEL: @aggr_loadstore_same_global
+; IR-NOT: addrspacecast
+; IR: icmp ult ptr addrspace(1) %a, %b
+; IR: load i8, ptr addrspace(1)
+; IR: store i8 {{.*}}, ptr addrspace(1)
+}
+
+define void @aggr_loadstore_noalias(ptr noalias %dst, ptr noalias %src) {
+ %v = load [128 x i8], ptr %src, align 1
+ store [128 x i8] %v, ptr %dst, align 1
+ ret void
+
+; noalias pointers can't overlap and BasicAA proves it (so this holds under opt
+; too). The copy is a plain forward loop carrying alias-scope/noalias metadata,
+; with no runtime direction check and no addrspacecast.
+; IR-LABEL: @aggr_loadstore_noalias
+; IR-NOT: memmove_bwd_loop
+; IR-NOT: addrspacecast
+; IR: load i8, ptr {{.*}}, !alias.scope
+; IR: store i8 {{.*}}, !noalias
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/201177
More information about the llvm-commits
mailing list