-.. _code_generator:
-
==========================================
The LLVM Target-Independent Code Generator
==========================================
.partial { background-color: #F88017 }
.yes { background-color: #0F0; }
.yes:before { content: "Y" }
+ .na { background-color: #6666FF; }
+ .na:before { content: "N/A" }
</style>
.. contents::
familiar with the `target description`_ and `machine code representation`_
classes. If you want to add a backend for a new target, you will need to
`implement the target description`_ classes for your new target and understand
-the `LLVM code representation <LangRef.html>`_. If you are interested in
+the :doc:`LLVM code representation <LangRef>`. If you are interested in
implementing a new `code generation algorithm`_, it should only depend on the
target-description and machine code representation classes, ensuring that it is
portable.
information (e.g., an ``add`` instruction is almost identical to a ``sub``
instruction). In order to allow the maximum amount of commonality to be
factored out, the LLVM code generator uses the
-`TableGen <TableGenFundamentals.html>`_ tool to describe big chunks of the
+:doc:`TableGen/index` tool to describe big chunks of the
target machine, which allows the use of domain-specific and target-specific
abstractions to reduce the amount of repetition.
pointers in the target, and whether the target is little-endian or
big-endian.
-.. _targetlowering:
+.. _TargetLowering:
The ``TargetLowering`` class
----------------------------
* various high-level characteristics, like whether it is profitable to turn
division by a constant into a multiplication sequence.
+.. _TargetRegisterInfo:
+
The ``TargetRegisterInfo`` class
--------------------------------
virtual registers with a physical register in the set.
The target-specific implementations of these classes is auto-generated from a
-`TableGen <TableGenFundamentals.html>`_ description of the register file.
+:doc:`TableGen/index` description of the register file.
.. _TargetInstrInfo:
-----------------------------
The ``TargetInstrInfo`` class is used to describe the machine instructions
-supported by the target. It is essentially an array of ``TargetInstrDescriptor``
-objects, each of which describes one instruction the target
-supports. Descriptors define things like the mnemonic for the opcode, the number
-of operands, the list of implicit register uses and defs, whether the
-instruction has certain target-independent properties (accesses memory, is
-commutable, etc), and holds any target-specific flags.
+supported by the target. Descriptions define things like the mnemonic for
+the opcode, the number of operands, the list of implicit register uses and defs,
+whether the instruction has certain target-independent properties (accesses
+memory, is commutable, etc), and holds any target-specific flags.
The ``TargetFrameInfo`` class
-----------------------------
.. code-block:: llvm
define i32 @test(i32 %X, i32 %Y) {
- %Z = udiv i32 %X, %Y
+ %Z = sdiv i32 %X, %Y
ret i32 %Z
}
-The X86 instruction selector produces this machine code for the ``div`` and
-``ret`` (use "``llc X.bc -march=x86 -print-machineinstrs``" to get this):
+The X86 instruction selector might produce this machine code for the ``div`` and
+``ret``:
.. code-block:: llvm
%EAX = mov %reg1026 ;; 32-bit return value goes in EAX
ret
-By the end of code generation, the register allocator has coalesced the
-registers and deleted the resultant identity moves producing the following
+By the end of code generation, the register allocator would coalesce the
+registers and delete the resultant identity moves producing the following
code:
.. code-block:: llvm
that prints out a directive for each method (e.g. ``EmitValue -> .byte``), but
MCObjectStreamer implements a full assembler.
+For target specific directives, the MCStreamer has a MCTargetStreamer instance.
+Each target that needs it defines a class that inherits from it and is a lot
+like MCStreamer itself: It has one method per directive and two classes that
+inherit from it, a target object streamer and a target asm streamer. The target
+asm streamer just prints it (``emitFnStart -> .fnstrart``), and the object
+streamer implement the assembler logic for it.
+
+To make llvm use these classes, the target initialization must call
+TargetRegistry::RegisterAsmStreamer and TargetRegistry::RegisterMCObjectStreamer
+passing callbacks that allocate the corresponding target streamer and pass it
+to createAsmStreamer or to the appropriate object streamer constructor.
+
The ``MCContext`` class
-----------------------
SREM or UREM operation. The `legalize types`_ and `legalize operations`_ phases
are responsible for turning an illegal DAG into a legal DAG.
+.. _SelectionDAG-Process:
+
SelectionDAG Instruction Selection Process
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
A target implementation tells the legalizer which types are supported (and which
register class to use for them) by calling the ``addRegisterClass`` method in
-its TargetLowering constructor.
+its ``TargetLowering`` constructor.
.. _legalize operations:
.. _Legalizer:
To perform this pattern match, the PowerPC backend includes the following
instruction definitions:
-::
+.. code-block:: text
+ :emphasize-lines: 4-5,9
def FMADDS : AForm_1<59, 29,
(ops F4RC:$FRT, F4RC:$FRA, F4RC:$FRC, F4RC:$FRB),
"fadds $FRT, $FRA, $FRB",
[(set F4RC:$FRT, (fadd F4RC:$FRA, F4RC:$FRB))]>;
-The portion of the instruction definition in bold indicates the pattern used to
-match the instruction. The DAG operators (like ``fmul``/``fadd``) are defined
-in the ``include/llvm/Target/TargetSelectionDAG.td`` file. " ``F4RC``" is the
-register class of the input and result values.
+The highlighted portion of the instruction definitions indicates the pattern
+used to match the instructions. The DAG operators (like ``fmul``/``fadd``)
+are defined in the ``include/llvm/Target/TargetSelectionDAG.td`` file.
+"``F4RC``" is the register class of the input and result values.
The TableGen DAG instruction selector generator reads the instruction patterns
in the ``.td`` file and automatically builds parts of the pattern matching code
are used to manipulate the input immediate (in this case, take the high or low
16-bits of the immediate).
+* When using the 'Pat' class to map a pattern to an instruction that has one
+ or more complex operands (like e.g. `X86 addressing mode`_), the pattern may
+ either specify the operand as a whole using a ``ComplexPattern``, or else it
+ may specify the components of the complex operand separately. The latter is
+ done e.g. for pre-increment instructions by the PowerPC back end:
+
+ ::
+
+ def STWU : DForm_1<37, (outs ptr_rc:$ea_res), (ins GPRC:$rS, memri:$dst),
+ "stwu $rS, $dst", LdStStoreUpd, []>,
+ RegConstraint<"$dst.reg = $ea_res">, NoEncode<"$ea_res">;
+
+ def : Pat<(pre_store GPRC:$rS, ptr_rc:$ptrreg, iaddroff:$ptroff),
+ (STWU GPRC:$rS, iaddroff:$ptroff, ptr_rc:$ptrreg)>;
+
+ Here, the pair of ``ptroff`` and ``ptrreg`` operands is matched onto the
+ complex operand ``dst`` of class ``memri`` in the ``STWU`` instruction.
+
* While the system does automate a lot, it still allows you to write custom C++
code to match special cases if there is something that is hard to
express.
===============================
Though you're probably reading this because you want to write or maintain a
-compiler backend, LLVM also fully supports building a native assemblers too.
+compiler backend, LLVM also fully supports building a native assembler.
We've tried hard to automate the generation of the assembler from the .td files
(in particular the instruction syntax and encodings), which means that a large
part of the manual and repetitive data entry can be factored and shared with the
specific to the code generator for a particular target. First we start with a
table that summarizes what features are supported by each target.
+.. _target-feature-matrix:
+
Target Feature Matrix
---------------------
:raw-html:`<table border="1" cellspacing="0">`
:raw-html:`<tr>`
:raw-html:`<th>Unknown</th>`
+:raw-html:`<th>Not Applicable</th>`
:raw-html:`<th>No support</th>`
:raw-html:`<th>Partial Support</th>`
:raw-html:`<th>Complete Support</th>`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td class="unknown"></td>`
+:raw-html:`<td class="na"></td>`
:raw-html:`<td class="no"></td>`
:raw-html:`<td class="partial"></td>`
:raw-html:`<td class="yes"></td>`
:raw-html:`<th>Feature</th>`
:raw-html:`<th>ARM</th>`
:raw-html:`<th>Hexagon</th>`
-:raw-html:`<th>MBlaze</th>`
:raw-html:`<th>MSP430</th>`
:raw-html:`<th>Mips</th>`
-:raw-html:`<th>PTX</th>`
+:raw-html:`<th>NVPTX</th>`
:raw-html:`<th>PowerPC</th>`
:raw-html:`<th>Sparc</th>`
+:raw-html:`<th>SystemZ</th>`
:raw-html:`<th>X86</th>`
:raw-html:`<th>XCore</th>`
:raw-html:`</tr>`
:raw-html:`<td><a href="#feat_reliable">is generally reliable</a></td>`
:raw-html:`<td class="yes"></td> <!-- ARM -->`
:raw-html:`<td class="yes"></td> <!-- Hexagon -->`
-:raw-html:`<td class="no"></td> <!-- MBlaze -->`
:raw-html:`<td class="unknown"></td> <!-- MSP430 -->`
:raw-html:`<td class="yes"></td> <!-- Mips -->`
-:raw-html:`<td class="no"></td> <!-- PTX -->`
+:raw-html:`<td class="yes"></td> <!-- NVPTX -->`
:raw-html:`<td class="yes"></td> <!-- PowerPC -->`
:raw-html:`<td class="yes"></td> <!-- Sparc -->`
+:raw-html:`<td class="yes"></td> <!-- SystemZ -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
-:raw-html:`<td class="unknown"></td> <!-- XCore -->`
+:raw-html:`<td class="yes"></td> <!-- XCore -->`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td><a href="#feat_asmparser">assembly parser</a></td>`
:raw-html:`<td class="no"></td> <!-- ARM -->`
:raw-html:`<td class="no"></td> <!-- Hexagon -->`
-:raw-html:`<td class="yes"></td> <!-- MBlaze -->`
:raw-html:`<td class="no"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
-:raw-html:`<td class="no"></td> <!-- PTX -->`
+:raw-html:`<td class="no"></td> <!-- NVPTX -->`
:raw-html:`<td class="no"></td> <!-- PowerPC -->`
:raw-html:`<td class="no"></td> <!-- Sparc -->`
+:raw-html:`<td class="yes"></td> <!-- SystemZ -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
:raw-html:`<td class="no"></td> <!-- XCore -->`
:raw-html:`</tr>`
:raw-html:`<td><a href="#feat_disassembler">disassembler</a></td>`
:raw-html:`<td class="yes"></td> <!-- ARM -->`
:raw-html:`<td class="no"></td> <!-- Hexagon -->`
-:raw-html:`<td class="yes"></td> <!-- MBlaze -->`
:raw-html:`<td class="no"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
-:raw-html:`<td class="no"></td> <!-- PTX -->`
+:raw-html:`<td class="na"></td> <!-- NVPTX -->`
:raw-html:`<td class="no"></td> <!-- PowerPC -->`
+:raw-html:`<td class="yes"></td> <!-- SystemZ -->`
:raw-html:`<td class="no"></td> <!-- Sparc -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
-:raw-html:`<td class="no"></td> <!-- XCore -->`
+:raw-html:`<td class="yes"></td> <!-- XCore -->`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td><a href="#feat_inlineasm">inline asm</a></td>`
:raw-html:`<td class="yes"></td> <!-- ARM -->`
:raw-html:`<td class="yes"></td> <!-- Hexagon -->`
-:raw-html:`<td class="yes"></td> <!-- MBlaze -->`
:raw-html:`<td class="unknown"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
-:raw-html:`<td class="unknown"></td> <!-- PTX -->`
+:raw-html:`<td class="yes"></td> <!-- NVPTX -->`
:raw-html:`<td class="yes"></td> <!-- PowerPC -->`
:raw-html:`<td class="unknown"></td> <!-- Sparc -->`
+:raw-html:`<td class="yes"></td> <!-- SystemZ -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
-:raw-html:`<td class="unknown"></td> <!-- XCore -->`
+:raw-html:`<td class="yes"></td> <!-- XCore -->`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td><a href="#feat_jit">jit</a></td>`
:raw-html:`<td class="partial"><a href="#feat_jit_arm">*</a></td> <!-- ARM -->`
:raw-html:`<td class="no"></td> <!-- Hexagon -->`
-:raw-html:`<td class="no"></td> <!-- MBlaze -->`
:raw-html:`<td class="unknown"></td> <!-- MSP430 -->`
:raw-html:`<td class="yes"></td> <!-- Mips -->`
-:raw-html:`<td class="unknown"></td> <!-- PTX -->`
+:raw-html:`<td class="na"></td> <!-- NVPTX -->`
:raw-html:`<td class="yes"></td> <!-- PowerPC -->`
:raw-html:`<td class="unknown"></td> <!-- Sparc -->`
+:raw-html:`<td class="yes"></td> <!-- SystemZ -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
-:raw-html:`<td class="unknown"></td> <!-- XCore -->`
+:raw-html:`<td class="no"></td> <!-- XCore -->`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td><a href="#feat_objectwrite">.o file writing</a></td>`
:raw-html:`<td class="no"></td> <!-- ARM -->`
:raw-html:`<td class="no"></td> <!-- Hexagon -->`
-:raw-html:`<td class="yes"></td> <!-- MBlaze -->`
:raw-html:`<td class="no"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
-:raw-html:`<td class="no"></td> <!-- PTX -->`
+:raw-html:`<td class="na"></td> <!-- NVPTX -->`
:raw-html:`<td class="no"></td> <!-- PowerPC -->`
:raw-html:`<td class="no"></td> <!-- Sparc -->`
+:raw-html:`<td class="yes"></td> <!-- SystemZ -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
:raw-html:`<td class="no"></td> <!-- XCore -->`
:raw-html:`</tr>`
:raw-html:`<td><a hr:raw-html:`ef="#feat_tailcall">tail calls</a></td>`
:raw-html:`<td class="yes"></td> <!-- ARM -->`
:raw-html:`<td class="yes"></td> <!-- Hexagon -->`
-:raw-html:`<td class="no"></td> <!-- MBlaze -->`
:raw-html:`<td class="unknown"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
-:raw-html:`<td class="unknown"></td> <!-- PTX -->`
+:raw-html:`<td class="no"></td> <!-- NVPTX -->`
:raw-html:`<td class="yes"></td> <!-- PowerPC -->`
:raw-html:`<td class="unknown"></td> <!-- Sparc -->`
+:raw-html:`<td class="no"></td> <!-- SystemZ -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
-:raw-html:`<td class="unknown"></td> <!-- XCore -->`
+:raw-html:`<td class="no"></td> <!-- XCore -->`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td><a href="#feat_segstacks">segmented stacks</a></td>`
:raw-html:`<td class="no"></td> <!-- ARM -->`
:raw-html:`<td class="no"></td> <!-- Hexagon -->`
-:raw-html:`<td class="no"></td> <!-- MBlaze -->`
:raw-html:`<td class="no"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
-:raw-html:`<td class="no"></td> <!-- PTX -->`
+:raw-html:`<td class="no"></td> <!-- NVPTX -->`
:raw-html:`<td class="no"></td> <!-- PowerPC -->`
:raw-html:`<td class="no"></td> <!-- Sparc -->`
+:raw-html:`<td class="no"></td> <!-- SystemZ -->`
:raw-html:`<td class="partial"><a href="#feat_segstacks_x86">*</a></td> <!-- X86 -->`
:raw-html:`<td class="no"></td> <!-- XCore -->`
:raw-html:`</tr>`
This box indicates whether the target supports guaranteed tail calls. These are
calls marked "`tail <LangRef.html#i_call>`_" and use the fastcc calling
-convention. Please see the `tail call section more more details`_.
+convention. Please see the `tail call section`_ for more details.
.. _feat_segstacks:
object files are not marked the way the gold linker expects, but simple Go
programs can be built by dragonegg.
-.. _tail call section more more details:
+.. _tail call section:
Tail call optimization
----------------------
TODO - More to come.
-The PTX backend
----------------
+The NVPTX backend
+-----------------
-The PTX code generator lives in the lib/Target/PTX directory. It is currently a
-work-in-progress, but already supports most of the code generation functionality
-needed to generate correct PTX kernels for CUDA devices.
+The NVPTX code generator under lib/Target/NVPTX is an open-source version of
+the NVIDIA NVPTX code generator for LLVM. It is contributed by NVIDIA and is
+a port of the code generator used in the CUDA compiler (nvcc). It targets the
+PTX 3.0/3.1 ISA and can target any compute capability greater than or equal to
+2.0 (Fermi).
-The code generator can target PTX 2.0+, and shader model 1.0+. The PTX ISA
-Reference Manual is used as the primary source of ISA information, though an
-effort is made to make the output of the code generator match the output of the
-NVidia nvcc compiler, whenever possible.
+This target is of production quality and should be completely compatible with
+the official NVIDIA toolchain.
Code Generator Options:
:raw-html:`<th>Description</th>`
:raw-html:`</tr>`
:raw-html:`<tr>`
-:raw-html:`<td>``double``</td>`
-:raw-html:`<td align="left">If enabled, the map_f64_to_f32 directive is disabled in the PTX output, allowing native double-precision arithmetic</td>`
+:raw-html:`<td>sm_20</td>`
+:raw-html:`<td align="left">Set shader model/compute capability to 2.0</td>`
:raw-html:`</tr>`
:raw-html:`<tr>`
-:raw-html:`<td>``no-fma``</td>`
-:raw-html:`<td align="left">Disable generation of Fused-Multiply Add instructions, which may be beneficial for some devices</td>`
+:raw-html:`<td>sm_21</td>`
+:raw-html:`<td align="left">Set shader model/compute capability to 2.1</td>`
:raw-html:`</tr>`
:raw-html:`<tr>`
-:raw-html:`<td>``smxy / computexy``</td>`
-:raw-html:`<td align="left">Set shader model/compute capability to x.y, e.g. sm20 or compute13</td>`
+:raw-html:`<td>sm_30</td>`
+:raw-html:`<td align="left">Set shader model/compute capability to 3.0</td>`
+:raw-html:`</tr>`
+:raw-html:`<tr>`
+:raw-html:`<td>sm_35</td>`
+:raw-html:`<td align="left">Set shader model/compute capability to 3.5</td>`
+:raw-html:`</tr>`
+:raw-html:`<tr>`
+:raw-html:`<td>ptx30</td>`
+:raw-html:`<td align="left">Target PTX 3.0</td>`
+:raw-html:`</tr>`
+:raw-html:`<tr>`
+:raw-html:`<td>ptx31</td>`
+:raw-html:`<td align="left">Target PTX 3.1</td>`
:raw-html:`</tr>`
:raw-html:`</table>`
-Working:
-
-* Arithmetic instruction selection (including combo FMA)
-
-* Bitwise instruction selection
-
-* Control-flow instruction selection
-
-* Function calls (only on SM 2.0+ and no return arguments)
-
-* Addresses spaces (0 = global, 1 = constant, 2 = local, 4 = shared)
-
-* Thread synchronization (bar.sync)
-
-* Special register reads ([N]TID, [N]CTAID, PMx, CLOCK, etc.)
-
-In Progress:
-
-* Robust call instruction selection
-
-* Stack frame allocation
-
-* Device-specific instruction scheduling optimizations
projects / oota-llvm.git / blobdiff