X-Git-Url: http://demsky.eecs.uci.edu/git/?a=blobdiff_plain;f=docs%2FCodeGenerator.rst;h=b5d41809744f9d310003fbb15a28181da93829a4;hb=624ec29474e1491a09bb99a783d1f27b7c4db841;hp=e2bbdb548cc07cffe51384e54bdc32d5d1d14d97;hpb=b129b9b617ea27df7ed176818ea9cb93d5d3b79d;p=oota-llvm.git diff --git a/docs/CodeGenerator.rst b/docs/CodeGenerator.rst index e2bbdb548cc..b5d41809744 100644 --- a/docs/CodeGenerator.rst +++ b/docs/CodeGenerator.rst @@ -1,5 +1,3 @@ -.. _code_generator: - ========================================== The LLVM Target-Independent Code Generator ========================================== @@ -17,6 +15,8 @@ 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" } .. contents:: @@ -250,6 +250,8 @@ operations. Among other things, this class indicates: * various high-level characteristics, like whether it is profitable to turn division by a constant into a multiplication sequence. +.. _TargetRegisterInfo: + The ``TargetRegisterInfo`` class -------------------------------- @@ -283,12 +285,10 @@ The ``TargetInstrInfo`` class ----------------------------- 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 ----------------------------- @@ -771,6 +771,8 @@ value of type i1, i8, i16, or i64 would be illegal, as would a DAG that uses a 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 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ @@ -1728,6 +1730,8 @@ This section of the document explains features or design decisions that are 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 --------------------- @@ -1742,12 +1746,14 @@ the key: :raw-html:`` :raw-html:`` :raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` @@ -1767,7 +1773,7 @@ Here is the table: :raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` @@ -1781,7 +1787,7 @@ Here is the table: :raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` @@ -1795,7 +1801,7 @@ Here is the table: :raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` @@ -1809,7 +1815,7 @@ Here is the table: :raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` @@ -1823,7 +1829,7 @@ Here is the table: :raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` @@ -1837,7 +1843,7 @@ Here is the table: :raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` @@ -1851,7 +1857,7 @@ Here is the table: :raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` @@ -1865,7 +1871,7 @@ Here is the table: :raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` @@ -1879,7 +1885,7 @@ Here is the table: :raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` :raw-html:`` @@ -2361,17 +2367,17 @@ Dynamic Allocation 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: @@ -2381,39 +2387,28 @@ Code Generator Options: :raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` -:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` -:raw-html:`` +:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`` -:raw-html:`` -:raw-html:`` +:raw-html:`` +:raw-html:`` :raw-html:`` :raw-html:`
UnknownNot ApplicableNo supportPartial SupportComplete Support
MBlazeMSP430MipsPTXNVPTXPowerPCSparcX86 * Description
``double``If enabled, the map_f64_to_f32 directive is disabled in the PTX output, allowing native double-precision arithmeticsm_20Set shader model/compute capability to 2.0
sm_21Set shader model/compute capability to 2.1
sm_30Set shader model/compute capability to 3.0
sm_35Set shader model/compute capability to 3.5
``no-fma``Disable generation of Fused-Multiply Add instructions, which may be beneficial for some devicesptx30Target PTX 3.0
``smxy / computexy``Set shader model/compute capability to x.y, e.g. sm20 or compute13ptx31Target PTX 3.1
` -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