Skip to content

Commit 2a5d458

Browse files
akshayrdeodharakshayrdeodhar
committed
cleanup
1 parent bbd7015 commit 2a5d458

File tree

15 files changed

+5041
-5073
lines changed

15 files changed

+5041
-5073
lines changed

‎llvm/include/llvm/CodeGen/TargetLowering.h‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2327,13 +2327,11 @@ class LLVM_ABI TargetLoweringBase {
23272327
/// @{
23282328
virtual Instruction *
23292329
emitLeadingFence(IRBuilderBase &Builder, Instruction *Inst,
2330-
AtomicOrdering Ord,
2331-
SyncScope::ID SSID = SyncScope::System) const;
2330+
AtomicOrdering Ord) const;
23322331

23332332
virtual Instruction *
23342333
emitTrailingFence(IRBuilderBase &Builder, Instruction *Inst,
2335-
AtomicOrdering Ord,
2336-
SyncScope::ID SSID = SyncScope::System) const;
2334+
AtomicOrdering Ord) const;
23372335
/// @}
23382336

23392337
// Emits code that executes when the comparison result in the ll/sc

‎llvm/lib/CodeGen/AtomicExpandPass.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -453,9 +453,9 @@ bool AtomicExpandImpl::bracketInstWithFences(Instruction *I,
453453
SyncScope::ID SSID) {
454454
ReplacementIRBuilder Builder(I, *DL);
455455

456-
auto LeadingFence = TLI->emitLeadingFence(Builder, I, Order, SSID);
456+
auto LeadingFence = TLI->emitLeadingFence(Builder, I, Order);
457457

458-
auto TrailingFence = TLI->emitTrailingFence(Builder, I, Order, SSID);
458+
auto TrailingFence = TLI->emitTrailingFence(Builder, I, Order);
459459
// We have a guard here because not every atomic operation generates a
460460
// trailing fence.
461461
if (TrailingFence)

‎llvm/lib/CodeGen/TargetLoweringBase.cpp‎

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -2383,20 +2383,18 @@ TargetLoweringBase::getAtomicMemOperandFlags(const Instruction &AI,
23832383

23842384
Instruction *TargetLoweringBase::emitLeadingFence(IRBuilderBase &Builder,
23852385
Instruction *Inst,
2386-
AtomicOrdering Ord,
2387-
SyncScope::ID SSID) const {
2386+
AtomicOrdering Ord) const {
23882387
if (isReleaseOrStronger(Ord) && Inst->hasAtomicStore())
2389-
return Builder.CreateFence(Ord, SSID);
2388+
return Builder.CreateFence(Ord);
23902389
else
23912390
return nullptr;
23922391
}
23932392

23942393
Instruction *TargetLoweringBase::emitTrailingFence(IRBuilderBase &Builder,
23952394
Instruction *Inst,
2396-
AtomicOrdering Ord,
2397-
SyncScope::ID SSID) const {
2395+
AtomicOrdering Ord) const {
23982396
if (isAcquireOrStronger(Ord))
2399-
return Builder.CreateFence(Ord, SSID);
2397+
return Builder.CreateFence(Ord);
24002398
else
24012399
return nullptr;
24022400
}

‎llvm/lib/Target/ARM/ARMISelLowering.cpp‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -21190,8 +21190,7 @@ Instruction *ARMTargetLowering::makeDMB(IRBuilderBase &Builder,
2119021190
// Based on http://www.cl.cam.ac.uk/~pes20/cpp/cpp0xmappings.html
2119121191
Instruction *ARMTargetLowering::emitLeadingFence(IRBuilderBase &Builder,
2119221192
Instruction *Inst,
21193-
AtomicOrdering Ord,
21194-
SyncScope::ID SSID) const {
21193+
AtomicOrdering Ord) const {
2119521194
switch (Ord) {
2119621195
case AtomicOrdering::NotAtomic:
2119721196
case AtomicOrdering::Unordered:
@@ -21216,8 +21215,7 @@ Instruction *ARMTargetLowering::emitLeadingFence(IRBuilderBase &Builder,
2121621215

2121721216
Instruction *ARMTargetLowering::emitTrailingFence(IRBuilderBase &Builder,
2121821217
Instruction *Inst,
21219-
AtomicOrdering Ord,
21220-
SyncScope::ID SSID) const {
21218+
AtomicOrdering Ord) const {
2122121219
switch (Ord) {
2122221220
case AtomicOrdering::NotAtomic:
2122321221
case AtomicOrdering::Unordered:

‎llvm/lib/Target/ARM/ARMISelLowering.h‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -675,11 +675,9 @@ class VectorType;
675675
emitAtomicCmpXchgNoStoreLLBalance(IRBuilderBase &Builder) const override;
676676

677677
Instruction *emitLeadingFence(
678-
IRBuilderBase &Builder, Instruction *Inst, AtomicOrdering Ord,
679-
SyncScope::ID SSID = SyncScope::ID SyncScope::System) const override;
678+
IRBuilderBase &Builder, Instruction *Inst, AtomicOrdering Ord) const override;
680679
Instruction *emitTrailingFence(
681-
IRBuilderBase &Builder, Instruction *Inst, AtomicOrdering Ord,
682-
SyncScope::ID SSID = SyncScope::ID SyncScope::System) const override;
680+
IRBuilderBase &Builder, Instruction *Inst, AtomicOrdering Ord) const override;
683681

684682
unsigned getMaxSupportedInterleaveFactor() const override;
685683

‎llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -6266,8 +6266,7 @@ AtomicOrdering NVPTXTargetLowering::atomicOperationOrderAfterFenceSplit(
62666266

62676267
Instruction *NVPTXTargetLowering::emitLeadingFence(IRBuilderBase &Builder,
62686268
Instruction *Inst,
6269-
AtomicOrdering Ord,
6270-
SyncScope::ID SSID) const {
6269+
AtomicOrdering Ord) const {
62716270
if (!isa<AtomicCmpXchgInst>(Inst))
62726271
return TargetLoweringBase::emitLeadingFence(Builder, Inst, Ord);
62736272

@@ -6285,8 +6284,7 @@ Instruction *NVPTXTargetLowering::emitLeadingFence(IRBuilderBase &Builder,
62856284

62866285
Instruction *NVPTXTargetLowering::emitTrailingFence(IRBuilderBase &Builder,
62876286
Instruction *Inst,
6288-
AtomicOrdering Ord,
6289-
SyncScope::ID SSID) const {
6287+
AtomicOrdering Ord) const {
62906288
// Specialize for cmpxchg
62916289
if (!isa<AtomicCmpXchgInst>(Inst))
62926290
return TargetLoweringBase::emitTrailingFence(Builder, Inst, Ord);

‎llvm/lib/Target/NVPTX/NVPTXInstrInfo.td‎

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1602,8 +1602,8 @@ def ADDR : Operand<pAny> {
16021602
let MIOperandInfo = (ops ADDR_base, i32imm);
16031603
}
16041604

1605-
def LdStCode : Operand<i32> {
1606-
let PrintMethod = "printLdStCode";
1605+
def AtomicCode : Operand<i32> {
1606+
let PrintMethod = "printAtomicCode";
16071607
}
16081608

16091609
def MmaCode : Operand<i32> {
@@ -1948,7 +1948,7 @@ defm ProxyRegB64 : ProxyRegInst<"b64", B64>;
19481948
class LD<NVPTXRegClass regclass>
19491949
: NVPTXInst<
19501950
(outs regclass:$dst),
1951-
(ins LdStCode:$sem, LdStCode:$scope, LdStCode:$addsp, LdStCode:$Sign,
1951+
(ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, AtomicCode:$Sign,
19521952
i32imm:$fromWidth, ADDR:$addr),
19531953
"ld${sem:sem}${scope:scope}${addsp:addsp}.${Sign:sign}$fromWidth "
19541954
"\t$dst, [$addr];", []>;
@@ -1964,7 +1964,7 @@ class ST<DAGOperand O>
19641964
: NVPTXInst<
19651965
(outs),
19661966
(ins O:$src,
1967-
LdStCode:$sem, LdStCode:$scope, LdStCode:$addsp, i32imm:$toWidth,
1967+
AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, i32imm:$toWidth,
19681968
ADDR:$addr),
19691969
"st${sem:sem}${scope:scope}${addsp:addsp}.b$toWidth"
19701970
" \t[$addr], $src;", []>;
@@ -1982,21 +1982,21 @@ let mayStore=1, hasSideEffects=0 in {
19821982
multiclass LD_VEC<NVPTXRegClass regclass, bit support_v8 = false> {
19831983
def _v2 : NVPTXInst<
19841984
(outs regclass:$dst1, regclass:$dst2),
1985-
(ins LdStCode:$sem, LdStCode:$scope, LdStCode:$addsp,
1986-
LdStCode:$Sign, i32imm:$fromWidth, ADDR:$addr),
1985+
(ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp,
1986+
AtomicCode:$Sign, i32imm:$fromWidth, ADDR:$addr),
19871987
"ld${sem:sem}${scope:scope}${addsp:addsp}.v2.${Sign:sign}$fromWidth "
19881988
"\t{{$dst1, $dst2}}, [$addr];", []>;
19891989
def _v4 : NVPTXInst<
19901990
(outs regclass:$dst1, regclass:$dst2, regclass:$dst3, regclass:$dst4),
1991-
(ins LdStCode:$sem, LdStCode:$scope, LdStCode:$addsp,
1992-
LdStCode:$Sign, i32imm:$fromWidth, ADDR:$addr),
1991+
(ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp,
1992+
AtomicCode:$Sign, i32imm:$fromWidth, ADDR:$addr),
19931993
"ld${sem:sem}${scope:scope}${addsp:addsp}.v4.${Sign:sign}$fromWidth "
19941994
"\t{{$dst1, $dst2, $dst3, $dst4}}, [$addr];", []>;
19951995
if support_v8 then
19961996
def _v8 : NVPTXInst<
19971997
(outs regclass:$dst1, regclass:$dst2, regclass:$dst3, regclass:$dst4,
19981998
regclass:$dst5, regclass:$dst6, regclass:$dst7, regclass:$dst8),
1999-
(ins LdStCode:$sem, LdStCode:$scope, LdStCode:$addsp, LdStCode:$Sign,
1999+
(ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, AtomicCode:$Sign,
20002000
i32imm:$fromWidth, ADDR:$addr),
20012001
"ld${sem:sem}${scope:scope}${addsp:addsp}.v8.${Sign:sign}$fromWidth "
20022002
"\t{{$dst1, $dst2, $dst3, $dst4, $dst5, $dst6, $dst7, $dst8}}, "
@@ -2013,14 +2013,14 @@ multiclass ST_VEC<DAGOperand O, bit support_v8 = false> {
20132013
def _v2 : NVPTXInst<
20142014
(outs),
20152015
(ins O:$src1, O:$src2,
2016-
LdStCode:$sem, LdStCode:$scope, LdStCode:$addsp, i32imm:$fromWidth,
2016+
AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, i32imm:$fromWidth,
20172017
ADDR:$addr),
20182018
"st${sem:sem}${scope:scope}${addsp:addsp}.v2.b$fromWidth "
20192019
"\t[$addr], {{$src1, $src2}};", []>;
20202020
def _v4 : NVPTXInst<
20212021
(outs),
20222022
(ins O:$src1, O:$src2, O:$src3, O:$src4,
2023-
LdStCode:$sem, LdStCode:$scope, LdStCode:$addsp, i32imm:$fromWidth,
2023+
AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, i32imm:$fromWidth,
20242024
ADDR:$addr),
20252025
"st${sem:sem}${scope:scope}${addsp:addsp}.v4.b$fromWidth "
20262026
"\t[$addr], {{$src1, $src2, $src3, $src4}};", []>;
@@ -2029,7 +2029,7 @@ multiclass ST_VEC<DAGOperand O, bit support_v8 = false> {
20292029
(outs),
20302030
(ins O:$src1, O:$src2, O:$src3, O:$src4,
20312031
O:$src5, O:$src6, O:$src7, O:$src8,
2032-
LdStCode:$sem, LdStCode:$scope, LdStCode:$addsp, i32imm:$fromWidth,
2032+
AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, i32imm:$fromWidth,
20332033
ADDR:$addr),
20342034
"st${sem:sem}${scope:scope}${addsp:addsp}.v8.b$fromWidth "
20352035
"\t[$addr], "

‎llvm/lib/Target/NVPTX/NVPTXIntrinsics.td‎

Lines changed: 9 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1882,22 +1882,22 @@ multiclass F_ATOMIC_2<RegTyInfo t, string sem_str, string as_str, string op_str,
18821882
}
18831883

18841884
multiclass F_ATOMIC_3<RegTyInfo t, string op_str> {
1885-
defvar asm_str = "atom${sem:sem}${scope:scope}${addsp:addsp}" # op_str
1885+
defvar asm_str = "atom${sem:sem}${scope:scope}${addsp:addsp}" # op_str # "\t$dst, [$addr], $b, $c;";
18861886

18871887
let mayLoad = 1, mayStore = 1, hasSideEffects = 1 in {
1888-
def _rr : BasicNVPTXInst<(outs t.RC:$dst),
1888+
def _rr : NVPTXInst<(outs t.RC:$dst),
18891889
(ins ADDR:$addr, t.RC:$b, t.RC:$c, AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp),
18901890
asm_str, []>;
18911891

1892-
def _ir : BasicNVPTXInst<(outs t.RC:$dst),
1892+
def _ir : NVPTXInst<(outs t.RC:$dst),
18931893
(ins ADDR:$addr, t.Imm:$b, t.RC:$c, AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp),
18941894
asm_str, []>;
18951895

1896-
def _ri : BasicNVPTXInst<(outs t.RC:$dst),
1896+
def _ri : NVPTXInst<(outs t.RC:$dst),
18971897
(ins ADDR:$addr, t.RC:$b, t.Imm:$c, AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp),
18981898
asm_str, []>;
18991899

1900-
def _ii : BasicNVPTXInst<(outs t.RC:$dst),
1900+
def _ii : NVPTXInst<(outs t.RC:$dst),
19011901
(ins ADDR:$addr, t.Imm:$b, t.Imm:$c, AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp),
19021902
asm_str, []>;
19031903
}
@@ -2030,10 +2030,6 @@ multiclass ATOM2S_impl<string OpStr, string IntTypeStr, string TypeStr,
20302030
multiclass F_ATOMIC_3_INTRINSIC_PATTERN<RegTyInfo t, string OpStr, string InstructionName, string IntTypeStr> {
20312031
foreach scope = ["cta", "sys"] in {
20322032
foreach space = ["gen"] in {
2033-
<<<<<<< HEAD
2034-
defm _#scope#space : ATOM3N_impl<OpStr, IntTypeStr, TypeStr, scope, space,
2035-
t, !listconcat(Preds, [hasAtomScope])>;
2036-
=======
20372033
defvar intrinsic = !cast<SDPatternOperator>("int_nvvm_atomic_" # OpStr # "_" # space # "_" # IntTypeStr # "_" # scope);
20382034
def : Pat<(t.Ty (intrinsic addr:$addr, t.Ty:$b, t.Ty:$c)),
20392035
(!cast<Instruction>(InstructionName # "_rr") ADDR:$addr, t.Ty:$b, t.Ty:$c, Ordering_not_atomic, !cast<PatLeaf>("Scope_" # scope), !cast<PatLeaf>("AddrSpace_" # space))>;
@@ -2046,7 +2042,6 @@ multiclass F_ATOMIC_3_INTRINSIC_PATTERN<RegTyInfo t, string OpStr, string Instru
20462042

20472043
def : Pat<(t.Ty (intrinsic addr:$addr, (t.Ty t.ImmNode:$b), (t.Ty t.ImmNode:$c))),
20482044
(!cast<Instruction>(InstructionName # "_ii") ADDR:$addr, (t.Ty t.ImmNode:$b), (t.Ty t.ImmNode:$c), Ordering_not_atomic, !cast<PatLeaf>("Scope_" # scope), !cast<PatLeaf>("AddrSpace_" # space))>;
2049-
>>>>>>> Replace old F_ATOMIC_3 completely with a single-opcode variation
20502045
}
20512046
}
20522047
}
@@ -2157,7 +2152,7 @@ def LDU_GLOBAL_v4i32 : VLDU_G_ELE_V4<"b32", B32>;
21572152
// during the lifetime of the kernel.
21582153

21592154
class LDG_G<NVPTXRegClass regclass>
2160-
: NVPTXInst<(outs regclass:$result), (ins LdStCode:$Sign, i32imm:$fromWidth, ADDR:$src),
2155+
: NVPTXInst<(outs regclass:$result), (ins AtomicCode:$Sign, i32imm:$fromWidth, ADDR:$src),
21612156
"ld.global.nc.${Sign:sign}$fromWidth \t$result, [$src];", []>;
21622157

21632158
def LD_GLOBAL_NC_i8 : LDG_G<B16>;
@@ -2170,19 +2165,19 @@ def LD_GLOBAL_NC_i64 : LDG_G<B64>;
21702165
// Elementized vector ldg
21712166
class VLDG_G_ELE_V2<NVPTXRegClass regclass> :
21722167
NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
2173-
(ins LdStCode:$Sign, i32imm:$fromWidth, ADDR:$src),
2168+
(ins AtomicCode:$Sign, i32imm:$fromWidth, ADDR:$src),
21742169
"ld.global.nc.v2.${Sign:sign}$fromWidth \t{{$dst1, $dst2}}, [$src];", []>;
21752170

21762171

21772172
class VLDG_G_ELE_V4<NVPTXRegClass regclass> :
21782173
NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3, regclass:$dst4),
2179-
(ins LdStCode:$Sign, i32imm:$fromWidth, ADDR:$src),
2174+
(ins AtomicCode:$Sign, i32imm:$fromWidth, ADDR:$src),
21802175
"ld.global.nc.v4.${Sign:sign}$fromWidth \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", []>;
21812176

21822177
class VLDG_G_ELE_V8<NVPTXRegClass regclass> :
21832178
NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3, regclass:$dst4,
21842179
regclass:$dst5, regclass:$dst6, regclass:$dst7, regclass:$dst8),
2185-
(ins LdStCode:$Sign, i32imm:$fromWidth, ADDR:$src),
2180+
(ins AtomicCode:$Sign, i32imm:$fromWidth, ADDR:$src),
21862181
"ld.global.nc.v8.${Sign:sign}$fromWidth \t{{$dst1, $dst2, $dst3, $dst4, $dst5, $dst6, $dst7, $dst8}}, [$src];", []>;
21872182

21882183
// FIXME: 8-bit LDG should be fixed once LDG/LDU nodes are made into proper loads.

‎llvm/lib/Target/PowerPC/PPCISelLowering.cpp‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -12820,8 +12820,7 @@ Value *PPCTargetLowering::emitStoreConditional(IRBuilderBase &Builder,
1282012820
// http://www.cl.cam.ac.uk/~pes20/cpp/cpp0xmappings.html
1282112821
Instruction *PPCTargetLowering::emitLeadingFence(IRBuilderBase &Builder,
1282212822
Instruction *Inst,
12823-
AtomicOrdering Ord,
12824-
SyncScope::ID SSID) const {
12823+
AtomicOrdering Ord) const {
1282512824
if (Ord == AtomicOrdering::SequentiallyConsistent)
1282612825
return callIntrinsic(Builder, Intrinsic::ppc_sync);
1282712826
if (isReleaseOrStronger(Ord))
@@ -12831,8 +12830,7 @@ Instruction *PPCTargetLowering::emitLeadingFence(IRBuilderBase &Builder,
1283112830

1283212831
Instruction *PPCTargetLowering::emitTrailingFence(IRBuilderBase &Builder,
1283312832
Instruction *Inst,
12834-
AtomicOrdering Ord,
12835-
SyncScope::ID SSID) const {
12833+
AtomicOrdering Ord) const {
1283612834
if (Inst->hasAtomicLoad() && isAcquireOrStronger(Ord)) {
1283712835
// See http://www.cl.cam.ac.uk/~pes20/cpp/cpp0xmappings.html and
1283812836
// http://www.rdrop.com/users/paulmck/scalability/paper/N2745r.2011.03.04a.html

‎llvm/lib/Target/PowerPC/PPCISelLowering.h‎

Lines changed: 0 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -927,7 +927,6 @@ namespace llvm {
927927
return true;
928928
}
929929

930-
<<<<<<< HEAD
931930
Value *emitLoadLinked(IRBuilderBase &Builder, Type *ValueTy, Value *Addr,
932931
AtomicOrdering Ord) const override;
933932

@@ -938,16 +937,6 @@ namespace llvm {
938937
AtomicOrdering Ord) const override;
939938
Instruction *emitTrailingFence(IRBuilderBase &Builder, Instruction *Inst,
940939
AtomicOrdering Ord) const override;
941-
=======
942-
Instruction *
943-
emitLeadingFence(IRBuilderBase &Builder, Instruction *Inst,
944-
AtomicOrdering Ord,
945-
SyncScope::ID SSID = SyncScope::System) const override;
946-
Instruction *
947-
emitTrailingFence(IRBuilderBase &Builder, Instruction *Inst,
948-
AtomicOrdering Ord,
949-
SyncScope::ID SSID = SyncScope::System) const override;
950-
>>>>>>> [NVPTX] Add syncscope support for cmpxchg
951940

952941
bool shouldInlineQuadwordAtomics() const;
953942

0 commit comments

Comments
 (0)