This is an archive of the discontinued LLVM Phabricator instance.

Differential D115597

Prevent store value forwarding to distinct addrspace load
ClosedPublic

Authored by dmitry on Dec 12 2021, 4:52 AM.

Details

Reviewers
hgreving
RKSimon
arsenm
spatel
Summary

DAGCombiner replaces (load const_addr1) directly chained with (store (val, const_addr2)) with val if address space stripped const_addr1 == const_addr2. The patch fixes the issue by checking address spaces as well.
However, it might makes sense to not to chain together side effects that belong to different address spaces in the first place and make SelectionDAG::root address space aware.

Diff Detail

Event Timeline

dmitry created this revision.Dec 12 2021, 4:52 AM
Herald added subscribers: asavonic, ecnelises, steven.zhang and 2 others. · View Herald TranscriptDec 12 2021, 4:52 AM
dmitry requested review of this revision.Dec 12 2021, 4:52 AM
Herald added a project: Restricted Project. · View Herald TranscriptDec 12 2021, 4:52 AM
Herald added subscribers: llvm-commits, wdng. · View Herald Transcript
dmitry edited the summary of this revision. (Show Details)Dec 12 2021, 4:54 AM
Harbormaster completed remote builds in B138857: Diff 393742.Dec 12 2021, 6:03 AM
arsenm accepted this revision.Dec 13 2021, 6:25 AM
This revision is now accepted and ready to land.Dec 13 2021, 6:25 AM
arsenm closed this revision.Dec 29 2022, 3:20 PM

Pushed as 0ec51a460a308f10a21bbaa9d96cd6c3e8e6d834

Herald added a project: Restricted Project. · View Herald TranscriptDec 29 2022, 3:20 PM
Herald added subscribers: mattd, gchakrabarti, arichardson. · View Herald Transcript

Revision Contents

PathSize
llvm/
lib/
CodeGen/
SelectionDAG/
2 lines
test/
CodeGen/
NVPTX/
10 lines

Diff 393742

llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 9,991 Lines • ▼ Show 20 Lines
} }
SDValue DAGCombiner::ForwardStoreValueToDirectLoad(LoadSDNode *LD) { SDValue DAGCombiner::ForwardStoreValueToDirectLoad(LoadSDNode *LD) {
if (OptLevel == CodeGenOpt::None || !LD->isSimple()) if (OptLevel == CodeGenOpt::None || !LD->isSimple())
return SDValue(); return SDValue();
SDValue Chain = LD->getOperand(0); SDValue Chain = LD->getOperand(0);
StoreSDNode *ST = dyn_cast<StoreSDNode>(Chain.getNode()); StoreSDNode *ST = dyn_cast<StoreSDNode>(Chain.getNode());
// TODO: Relax this restriction for unordered atomics (see D66309) // TODO: Relax this restriction for unordered atomics (see D66309)
if (!ST || !ST->isSimple()) if (!ST || !ST->isSimple() || ST->getAddressSpace() != LD->getAddressSpace())
return SDValue(); return SDValue();
EVT LDType = LD->getValueType(0); EVT LDType = LD->getValueType(0);
EVT LDMemType = LD->getMemoryVT(); EVT LDMemType = LD->getMemoryVT();
EVT STMemType = ST->getMemoryVT(); EVT STMemType = ST->getMemoryVT();
EVT STType = ST->getValue().getValueType(); EVT STType = ST->getValue().getValueType();
// There are two cases to consider here: // There are two cases to consider here:
▲ Show 20 Lines • Show All 7,903 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/chain-different-as.ll

  • This file was added.
; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s
define i64 @test() nounwind readnone {
%addr0 = inttoptr i64 1 to i64*
; CHECK: mov.u64 %[[reg:rd[0-9]+]], 1
; CHECK: ld.global.u64 %{{rd[0-9]+}}, [%[[reg]]]
%addr1 = inttoptr i64 1 to i64 addrspace(1)*
store i64 42, i64* %addr0
%res = load i64, i64 addrspace(1)* %addr1
ret i64 %res
}