Skip to content

Commit de11258

Browse files
AlexMacleanAlexMaclean
committed
nfc cleanup 2
1 parent 80733cb commit de11258

File tree

2 files changed

+268
-484
lines changed

2 files changed

+268
-484
lines changed

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

Lines changed: 43 additions & 51 deletions
Original file line numberDiff line numberDiff line change
@@ -268,7 +268,7 @@ multiclass I3Inst<string op_str, SDPatternOperator op_node, RegTyInfo t,
268268
// The instructions are named "<OpcStr><Width>" (e.g. "add.s64").
269269
multiclass I3<string op_str, SDPatternOperator op_node, bit commutative> {
270270
foreach t = [I16RT, I32RT, I64RT] in
271-
defm t.Ty# : I3Inst<op_str # t.Size, op_node, t, commutative>;
271+
defm t.Size# : I3Inst<op_str # t.Size, op_node, t, commutative>;
272272
}
273273

274274
class I16x2<string OpcStr, SDNode OpNode> :
@@ -757,8 +757,8 @@ defm SUBCCC : ADD_SUB_INT_CARRY<"subc.cc", sube, commutative = false>;
757757

758758
defm MULT : I3<"mul.lo.s", mul, commutative = true>;
759759

760-
defm MULTHS : I3<"mul.hi.s", mulhs, commutative = true>;
761-
defm MULTHU : I3<"mul.hi.u", mulhu, commutative = true>;
760+
defm MUL_HI_S : I3<"mul.hi.s", mulhs, commutative = true>;
761+
defm MUL_HI_U : I3<"mul.hi.u", mulhu, commutative = true>;
762762

763763
defm SDIV : I3<"div.s", sdiv, commutative = false>;
764764
defm UDIV : I3<"div.u", udiv, commutative = false>;
@@ -977,7 +977,7 @@ def fdiv_approx : PatFrag<(ops node:$a, node:$b),
977977
}]>;
978978

979979

980-
def FRCP32_approx_r :
980+
def RCP_APPROX_F32_r :
981981
BasicFlagsNVPTXInst<(outs B32:$dst),
982982
(ins B32:$b), (ins FTZFlag:$ftz),
983983
"rcp.approx$ftz.f32",
@@ -986,12 +986,12 @@ def FRCP32_approx_r :
986986
//
987987
// F32 Approximate division
988988
//
989-
def FDIV32_approx_rr :
989+
def DIV_APPROX_F32_rr :
990990
BasicFlagsNVPTXInst<(outs B32:$dst),
991991
(ins B32:$a, B32:$b), (ins FTZFlag:$ftz),
992992
"div.approx$ftz.f32",
993993
[(set f32:$dst, (fdiv_approx f32:$a, f32:$b))]>;
994-
def FDIV32_approx_ri :
994+
def DIV_APPROX_F32_ri :
995995
BasicFlagsNVPTXInst<(outs B32:$dst),
996996
(ins B32:$a, f32imm:$b), (ins FTZFlag:$ftz),
997997
"div.approx$ftz.f32",
@@ -1009,7 +1009,7 @@ def fdiv_full : PatFrag<(ops node:$a, node:$b),
10091009

10101010

10111011
def : Pat<(fdiv_full f32imm_1, f32:$b),
1012-
(FRCP32_approx_r $b)>;
1012+
(RCP_APPROX_F32_r $b)>;
10131013

10141014
//
10151015
// F32 Semi-accurate division
@@ -1475,9 +1475,9 @@ def MmaCode : Operand<i32> {
14751475
// Get pointer to local stack.
14761476
let hasSideEffects = false in {
14771477
def MOV_DEPOT_ADDR : NVPTXInst<(outs B32:$d), (ins i32imm:$num),
1478-
"mov.b32 \t$d, __local_depot$num;", []>;
1478+
"mov.b32 \t$d, __local_depot$num;">;
14791479
def MOV_DEPOT_ADDR_64 : NVPTXInst<(outs B64:$d), (ins i32imm:$num),
1480-
"mov.b64 \t$d, __local_depot$num;", []>;
1480+
"mov.b64 \t$d, __local_depot$num;">;
14811481
}
14821482

14831483

@@ -1533,9 +1533,9 @@ def : Pat<(i64 externalsym:$dst), (MOV_B64_i (to_texternsym $dst))>;
15331533

15341534
//---- Copy Frame Index ----
15351535
def LEA_ADDRi : NVPTXInst<(outs B32:$dst), (ins ADDR:$addr),
1536-
"add.u32 \t$dst, ${addr:add};", []>;
1536+
"add.u32 \t$dst, ${addr:add};">;
15371537
def LEA_ADDRi64 : NVPTXInst<(outs B64:$dst), (ins ADDR:$addr),
1538-
"add.u64 \t$dst, ${addr:add};", []>;
1538+
"add.u64 \t$dst, ${addr:add};">;
15391539

15401540
def : Pat<(i32 frameindex:$fi), (LEA_ADDRi (to_tframeindex $fi), 0)>;
15411541
def : Pat<(i64 frameindex:$fi), (LEA_ADDRi64 (to_tframeindex $fi), 0)>;
@@ -1612,12 +1612,12 @@ foreach is_convergent = [0, 1] in {
16121612
NVPTXInst<(outs),
16131613
(ins ADDR_base:$addr, CallOperand:$rets, CallOperand:$params,
16141614
i32imm:$proto),
1615-
"call${rets:RetList} $addr, (${params:ParamList}), prototype_$proto;", []>;
1615+
"call${rets:RetList} $addr, (${params:ParamList}), prototype_$proto;">;
16161616

16171617
def CALL_UNI # convergent_suffix :
16181618
NVPTXInst<(outs),
16191619
(ins ADDR_base:$addr, CallOperand:$rets, CallOperand:$params),
1620-
"call.uni${rets:RetList} $addr, (${params:ParamList});", []>;
1620+
"call.uni${rets:RetList} $addr, (${params:ParamList});">;
16211621
}
16221622

16231623
defvar call_inst = !cast<NVPTXInst>("CALL" # convergent_suffix);
@@ -1633,10 +1633,10 @@ foreach is_convergent = [0, 1] in {
16331633

16341634
def DECLARE_PARAM_array :
16351635
NVPTXInst<(outs), (ins i32imm:$a, i32imm:$align, i32imm:$size),
1636-
".param .align $align .b8 \t$a[$size];", []>;
1636+
".param .align $align .b8 \t$a[$size];">;
16371637
def DECLARE_PARAM_scalar :
16381638
NVPTXInst<(outs), (ins i32imm:$a, i32imm:$size),
1639-
".param .b$size \t$a;", []>;
1639+
".param .b$size \t$a;">;
16401640

16411641
def : Pat<(declare_array_param externalsym:$a, imm:$align, imm:$size),
16421642
(DECLARE_PARAM_array (to_texternsym $a), imm:$align, imm:$size)>;
@@ -1709,7 +1709,7 @@ class LD<NVPTXRegClass regclass>
17091709
(ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, AtomicCode:$Sign,
17101710
i32imm:$fromWidth, ADDR:$addr),
17111711
"ld${sem:sem}${scope:scope}${addsp:addsp}.${Sign:sign}$fromWidth "
1712-
"\t$dst, [$addr];", []>;
1712+
"\t$dst, [$addr];">;
17131713

17141714
let mayLoad=1, hasSideEffects=0 in {
17151715
def LD_i16 : LD<B16>;
@@ -1724,7 +1724,7 @@ class ST<DAGOperand O>
17241724
AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, i32imm:$toWidth,
17251725
ADDR:$addr),
17261726
"st${sem:sem}${scope:scope}${addsp:addsp}.b$toWidth"
1727-
" \t[$addr], $src;", []>;
1727+
" \t[$addr], $src;">;
17281728

17291729
let mayStore=1, hasSideEffects=0 in {
17301730
def ST_i16 : ST<RI16>;
@@ -1741,13 +1741,13 @@ multiclass LD_VEC<NVPTXRegClass regclass, bit support_v8 = false> {
17411741
(ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp,
17421742
AtomicCode:$Sign, i32imm:$fromWidth, ADDR:$addr),
17431743
"ld${sem:sem}${scope:scope}${addsp:addsp}.v2.${Sign:sign}$fromWidth "
1744-
"\t{{$dst1, $dst2}}, [$addr];", []>;
1744+
"\t{{$dst1, $dst2}}, [$addr];">;
17451745
def _v4 : NVPTXInst<
17461746
(outs regclass:$dst1, regclass:$dst2, regclass:$dst3, regclass:$dst4),
17471747
(ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp,
17481748
AtomicCode:$Sign, i32imm:$fromWidth, ADDR:$addr),
17491749
"ld${sem:sem}${scope:scope}${addsp:addsp}.v4.${Sign:sign}$fromWidth "
1750-
"\t{{$dst1, $dst2, $dst3, $dst4}}, [$addr];", []>;
1750+
"\t{{$dst1, $dst2, $dst3, $dst4}}, [$addr];">;
17511751
if support_v8 then
17521752
def _v8 : NVPTXInst<
17531753
(outs regclass:$dst1, regclass:$dst2, regclass:$dst3, regclass:$dst4,
@@ -1756,7 +1756,7 @@ multiclass LD_VEC<NVPTXRegClass regclass, bit support_v8 = false> {
17561756
i32imm:$fromWidth, ADDR:$addr),
17571757
"ld${sem:sem}${scope:scope}${addsp:addsp}.v8.${Sign:sign}$fromWidth "
17581758
"\t{{$dst1, $dst2, $dst3, $dst4, $dst5, $dst6, $dst7, $dst8}}, "
1759-
"[$addr];", []>;
1759+
"[$addr];">;
17601760
}
17611761
let mayLoad=1, hasSideEffects=0 in {
17621762
defm LDV_i16 : LD_VEC<B16>;
@@ -1771,14 +1771,14 @@ multiclass ST_VEC<DAGOperand O, bit support_v8 = false> {
17711771
AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, i32imm:$fromWidth,
17721772
ADDR:$addr),
17731773
"st${sem:sem}${scope:scope}${addsp:addsp}.v2.b$fromWidth "
1774-
"\t[$addr], {{$src1, $src2}};", []>;
1774+
"\t[$addr], {{$src1, $src2}};">;
17751775
def _v4 : NVPTXInst<
17761776
(outs),
17771777
(ins O:$src1, O:$src2, O:$src3, O:$src4,
17781778
AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, i32imm:$fromWidth,
17791779
ADDR:$addr),
17801780
"st${sem:sem}${scope:scope}${addsp:addsp}.v4.b$fromWidth "
1781-
"\t[$addr], {{$src1, $src2, $src3, $src4}};", []>;
1781+
"\t[$addr], {{$src1, $src2, $src3, $src4}};">;
17821782
if support_v8 then
17831783
def _v8 : NVPTXInst<
17841784
(outs),
@@ -1788,7 +1788,7 @@ multiclass ST_VEC<DAGOperand O, bit support_v8 = false> {
17881788
ADDR:$addr),
17891789
"st${sem:sem}${scope:scope}${addsp:addsp}.v8.b$fromWidth "
17901790
"\t[$addr], "
1791-
"{{$src1, $src2, $src3, $src4, $src5, $src6, $src7, $src8}};", []>;
1791+
"{{$src1, $src2, $src3, $src4, $src5, $src6, $src7, $src8}};">;
17921792
}
17931793

17941794
let mayStore=1, hasSideEffects=0 in {
@@ -1983,60 +1983,52 @@ let hasSideEffects = false in {
19831983
def V4I16toI64 : NVPTXInst<(outs B64:$d),
19841984
(ins B16:$s1, B16:$s2,
19851985
B16:$s3, B16:$s4),
1986-
"mov.b64 \t$d, {{$s1, $s2, $s3, $s4}};", []>;
1986+
"mov.b64 \t$d, {{$s1, $s2, $s3, $s4}};">;
19871987
def V2I16toI32 : NVPTXInst<(outs B32:$d),
19881988
(ins B16:$s1, B16:$s2),
1989-
"mov.b32 \t$d, {{$s1, $s2}};", []>;
1989+
"mov.b32 \t$d, {{$s1, $s2}};">;
19901990
def V2I32toI64 : NVPTXInst<(outs B64:$d),
19911991
(ins B32:$s1, B32:$s2),
1992-
"mov.b64 \t$d, {{$s1, $s2}};", []>;
1992+
"mov.b64 \t$d, {{$s1, $s2}};">;
19931993
def V2I64toI128 : NVPTXInst<(outs B128:$d),
19941994
(ins B64:$s1, B64:$s2),
1995-
"mov.b128 \t$d, {{$s1, $s2}};", []>;
1995+
"mov.b128 \t$d, {{$s1, $s2}};">;
19961996

19971997
// unpack a larger int register to a set of smaller int registers
19981998
def I64toV4I16 : NVPTXInst<(outs B16:$d1, B16:$d2,
19991999
B16:$d3, B16:$d4),
20002000
(ins B64:$s),
2001-
"mov.b64 \t{{$d1, $d2, $d3, $d4}}, $s;", []>;
2001+
"mov.b64 \t{{$d1, $d2, $d3, $d4}}, $s;">;
20022002
def I32toV2I16 : NVPTXInst<(outs B16:$d1, B16:$d2),
20032003
(ins B32:$s),
2004-
"mov.b32 \t{{$d1, $d2}}, $s;", []>;
2004+
"mov.b32 \t{{$d1, $d2}}, $s;">;
20052005
def I64toV2I32 : NVPTXInst<(outs B32:$d1, B32:$d2),
20062006
(ins B64:$s),
2007-
"mov.b64 \t{{$d1, $d2}}, $s;", []>;
2007+
"mov.b64 \t{{$d1, $d2}}, $s;">;
20082008
def I128toV2I64: NVPTXInst<(outs B64:$d1, B64:$d2),
20092009
(ins B128:$s),
2010-
"mov.b128 \t{{$d1, $d2}}, $s;", []>;
2010+
"mov.b128 \t{{$d1, $d2}}, $s;">;
20112011

2012-
def I32toI16H : NVPTXInst<(outs B16:$high),
2013-
(ins B32:$s),
2014-
"{{ .reg .b16 tmp; mov.b32 {tmp, $high}, $s; }}",
2015-
[]>;
2016-
def I32toI16L : NVPTXInst<(outs B16:$low),
2017-
(ins B32:$s),
2018-
"{{ .reg .b16 tmp; mov.b32 {$low, tmp}, $s; }}",
2019-
[]>;
2020-
def I64toI32H : NVPTXInst<(outs B32:$high),
2021-
(ins B64:$s),
2022-
"{{ .reg .b32 tmp; mov.b64 {tmp, $high}, $s; }}",
2023-
[]>;
2024-
def I64toI32L : NVPTXInst<(outs B32:$low),
2025-
(ins B64:$s),
2026-
"{{ .reg .b32 tmp; mov.b64 {$low, tmp}, $s; }}",
2027-
[]>;
2012+
def I32toI16H : NVPTXInst<(outs B16:$high), (ins B32:$s),
2013+
"{{ .reg .b16 tmp; mov.b32 {tmp, $high}, $s; }}">;
2014+
def I32toI16L : NVPTXInst<(outs B16:$low), (ins B32:$s),
2015+
"{{ .reg .b16 tmp; mov.b32 {$low, tmp}, $s; }}">;
2016+
def I64toI32H : NVPTXInst<(outs B32:$high), (ins B64:$s),
2017+
"{{ .reg .b32 tmp; mov.b64 {tmp, $high}, $s; }}">;
2018+
def I64toI32L : NVPTXInst<(outs B32:$low), (ins B64:$s),
2019+
"{{ .reg .b32 tmp; mov.b64 {$low, tmp}, $s; }}">;
20282020

20292021
// PTX 7.1 lets you avoid a temp register and just use _ as a "sink" for the
20302022
// unused high/low part.
20312023
let Predicates = [hasPTX<71>] in {
20322024
def I32toI16H_Sink : NVPTXInst<(outs B16:$high), (ins B32:$s),
2033-
"mov.b32 \t{{_, $high}}, $s;", []>;
2025+
"mov.b32 \t{{_, $high}}, $s;">;
20342026
def I32toI16L_Sink : NVPTXInst<(outs B16:$low), (ins B32:$s),
2035-
"mov.b32 \t{{$low, _}}, $s;", []>;
2027+
"mov.b32 \t{{$low, _}}, $s;">;
20362028
def I64toI32H_Sink : NVPTXInst<(outs B32:$high), (ins B64:$s),
2037-
"mov.b64 \t{{_, $high}}, $s;", []>;
2029+
"mov.b64 \t{{_, $high}}, $s;">;
20382030
def I64toI32L_Sink : NVPTXInst<(outs B32:$low), (ins B64:$s),
2039-
"mov.b64 \t{{$low, _}}, $s;", []>;
2031+
"mov.b64 \t{{$low, _}}, $s;">;
20402032
}
20412033
}
20422034

0 commit comments

Comments
 (0)