This is an archive of the discontinued LLVM Phabricator instance.

Differential D17432

[NVPTX] Nix hack used to emit '{' and '}' for NVPTX calls.
ClosedPublic

Authored by jlebar on Feb 18 2016, 6:07 PM.

Details

Reviewers
jholewinski
rnk
Commits
rG93e7a9b91c93: [NVPTX] Nix hack used to emit '{' and '}' for NVPTX calls.
rL262372: [NVPTX] Nix hack used to emit '{' and '}' for NVPTX calls.
Summary

Tablegen understands backslash as an escape char; that's sufficient.

Diff Detail

Repository
rL LLVM

Event Timeline

jlebar updated this revision to Diff 48443.Feb 18 2016, 6:07 PM
jlebar retitled this revision from to [NVPTX] Nix hack used to emit '{' and '}' for NVPTX calls..
jlebar updated this object.
jlebar added a reviewer: jholewinski.
jlebar added subscribers: tra, llvm-commits.
Herald added a subscriber: jholewinski. · View Herald TranscriptFeb 18 2016, 6:07 PM
jlebar added a reviewer: rnk.Feb 29 2016, 5:39 PM
rnk accepted this revision.Feb 29 2016, 6:00 PM
rnk edited edge metadata.

lgtm

This revision is now accepted and ready to land.Feb 29 2016, 6:00 PM
Closed by commit rL262372: [NVPTX] Nix hack used to emit '{' and '}' for NVPTX calls. (authored by jlebar). · Explain WhyMar 1 2016, 11:28 AM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
llvm/
trunk/
lib/
Target/
NVPTX/
12 lines

Diff 49519

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

Show First 20 Lines • Show All 2,676 Lines • ▼ Show 20 Lines
def : Pat<(call tglobaladdr:$dst), (CALL tglobaladdr:$dst)>; def : Pat<(call tglobaladdr:$dst), (CALL tglobaladdr:$dst)>;
def : Pat<(call texternalsym:$dst), (CALL texternalsym:$dst)>; def : Pat<(call texternalsym:$dst), (CALL texternalsym:$dst)>;
// Pseudo instructions. // Pseudo instructions.
class Pseudo<dag outs, dag ins, string asmstr, list<dag> pattern> class Pseudo<dag outs, dag ins, string asmstr, list<dag> pattern>
: NVPTXInst<outs, ins, asmstr, pattern>; : NVPTXInst<outs, ins, asmstr, pattern>;
// @TODO: We use some tricks here to emit curly braces. Can we clean this up
// a bit without TableGen modifications?
def Callseq_Start : def Callseq_Start :
NVPTXInst<(outs), (ins i32imm:$amt), NVPTXInst<(outs), (ins i32imm:$amt),
"// Callseq Start $amt\n" "\\{ // callseq $amt\n"
"\t{{\n" "\t.reg .b32 temp_param_reg;",
"\t.reg .b32 temp_param_reg;\n"
"\t// <end>}}",
[(callseq_start timm:$amt)]>; [(callseq_start timm:$amt)]>;
def Callseq_End : def Callseq_End :
NVPTXInst<(outs), (ins i32imm:$amt1, i32imm:$amt2), NVPTXInst<(outs), (ins i32imm:$amt1, i32imm:$amt2),
"\n" "\\} // callseq $amt1",
"\t//{{\n"
"\t}}// Callseq End $amt1",
[(callseq_end timm:$amt1, timm:$amt2)]>; [(callseq_end timm:$amt1, timm:$amt2)]>;
// trap instruction // trap instruction
def trapinst : NVPTXInst<(outs), (ins), "trap;", [(trap)]>; def trapinst : NVPTXInst<(outs), (ins), "trap;", [(trap)]>;
// Call prototype wrapper // Call prototype wrapper
def SDTCallPrototype : SDTypeProfile<0, 1, [SDTCisInt<0>]>; def SDTCallPrototype : SDTypeProfile<0, 1, [SDTCisInt<0>]>;
def CallPrototype : def CallPrototype :
Show All 21 Lines