<li><a href="#ppc_prolog">Prolog/Epilog</a></li>
<li><a href="#ppc_dynamic">Dynamic Allocation</a></li>
</ul></li>
+ <li><a href="#ptx">The PTX backend</a></li>
</ul></li>
</ol>
different register allocators:</p>
<ul>
- <li><i>Linear Scan</i> — <i>The default allocator</i>. This is the
- well-know linear scan register allocator. Whereas the
- <i>Simple</i> and <i>Local</i> algorithms use a direct mapping
- implementation technique, the <i>Linear Scan</i> implementation
- uses a spiller in order to place load and stores.</li>
-
<li><i>Fast</i> — This register allocator is the default for debug
builds. It allocates registers on a basic block level, attempting to keep
values in registers and reusing registers as appropriate.</li>
+ <li><i>Basic</i> — This is an incremental approach to register
+ allocation. Live ranges are assigned to registers one at a time in
+ an order that is driven by heuristics. Since code can be rewritten
+ on-the-fly during allocation, this framework allows interesting
+ allocators to be developed as extensions. It is not itself a
+ production register allocator but is a potentially useful
+ stand-alone mode for triaging bugs and as a performance baseline.
+
+ <li><i>Greedy</i> — <i>The default allocator</i>. This is a
+ highly tuned implementation of the <i>Basic</i> allocator that
+ incorporates global live range splitting. This allocator works hard
+ to minimize the cost of spill code.
+
<li><i>PBQP</i> — A Partitioned Boolean Quadratic Programming (PBQP)
based register allocator. This allocator works by constructing a PBQP
problem representing the register allocation problem under consideration,
solving this using a PBQP solver, and mapping the solution back to a
register assignment.</li>
-
</ul>
<p>The type of register allocator used in <tt>llc</tt> can be chosen with the
<div>
-<p>Unwinding out of a function is done virually via DWARF encodings. These
- encodings exist in two forms: a Common Information Entry (CIE) and a Frame
- Description Entry (FDE). These two tables contain the information necessary
- for the unwinder to restore the state of the computer to before the function
- was called. However, the tables themselves are rather large. LLVM can use a
- "compact unwind" encoding to represent the virtual unwinding.</p>
+<p>Throwing an exception requires <em>unwinding</em> out of a function. The
+ information on how to unwind a given function is traditionally expressed in
+ DWARF unwind (a.k.a. frame) info. But that format was originally developed
+ for debuggers to backtrace, and each Frame Description Entry (FDE) requires
+ ~20-30 bytes per function. There is also the cost of mapping from an address
+ in a function to the corresponding FDE at runtime. An alternative unwind
+ encoding is called <em>compact unwind</em> and requires just 4-bytes per
+ function.</p>
<p>The compact unwind encoding is a 32-bit value, which is encoded in an
architecture-specific way. It specifies which registers to restore and from
- where, and how to unwind out of the funciton. When the linker creates a final
+ where, and how to unwind out of the function. When the linker creates a final
linked image, it will create a <code>__TEXT,__unwind_info</code>
section. This section is a small and fast way for the runtime to access
unwind info for any given function. If we emit compact unwind info for the
<p>For X86, there are three modes for the compact unwind encoding:</p>
-<ul>
+<dl>
<dt><i>Function with a Frame Pointer (<code>EBP</code> or <code>RBP</code>)</i></dt>
<dd><p><code>EBP/RBP</code>-based frame, where <code>EBP/RBP</code> is pushed
onto the stack immediately after the return address,
more into the PC. All non-volatile registers that need to be restored must
have been saved in a small range on the stack that
starts <code>EBP-4</code> to <code>EBP-1020</code> (<code>RBP-8</code>
- to <code>RBP-1020</code>). The offset (divided by 4) is encoded in bits
- 16-23 (mask: <code>0x00FF0000</code>). The registers saved are encoded in
- bits 0-14 (mask: <code>0x00007FFF</code>) as five 3-bit entries from the
- following table:</p>
+ to <code>RBP-1020</code>). The offset (divided by 4 in 32-bit mode and 8
+ in 64-bit mode) is encoded in bits 16-23 (mask: <code>0x00FF0000</code>).
+ The registers saved are encoded in bits 0-14
+ (mask: <code>0x00007FFF</code>) as five 3-bit entries from the following
+ table:</p>
<table border="1" cellspacing="0">
<tr>
<th>Compact Number</th>
to the <code>ESP/RSP</code>. Then the return is done by popping the stack
into the PC. All non-volatile registers that need to be restored must have
been saved on the stack immediately after the return address. The stack
- size (divided by 4) is encoded in bits 16-23
- (mask: <code>0x00FF0000</code>). There is a maximum stack size of 1024
- bytes. The number of registers saved is encoded in bits 9-12
- (mask: <code>0x00001C00</code>). Bits 0-9 (mask:
- <code>0x000003FF</code>) contain which registers were saved and their
- order. (See the <code>encodeCompactUnwindRegistersWithoutFrame()</code>
- function in <code>lib/Target/X86FrameLowering.cpp</code> for the encoding
+ size (divided by 4 in 32-bit mode and 8 in 64-bit mode) is encoded in bits
+ 16-23 (mask: <code>0x00FF0000</code>). There is a maximum stack size of
+ 1024 bytes in 32-bit mode and 2048 in 64-bit mode. The number of registers
+ saved is encoded in bits 9-12 (mask: <code>0x00001C00</code>). Bits 0-9
+ (mask: <code>0x000003FF</code>) contain which registers were saved and
+ their order. (See
+ the <code>encodeCompactUnwindRegistersWithoutFrame()</code> function
+ in <code>lib/Target/X86FrameLowering.cpp</code> for the encoding
algorithm.)</p></dd>
<dt><i>Frameless with a Large Constant Stack Size (<code>EBP</code>
$nnnnnn, %esp</code>" in its prolog. The compact encoding contains the
offset to the <code>$nnnnnn</code> value in the function in bits 9-12
(mask: <code>0x00001C00</code>).</p></dd>
-</ul>
+</dl>
</div>
</div>
+</div>
+
+<!-- ======================================================================= -->
+<h3>
+ <a name="ptx">The PTX backend</a>
+</h3>
+
+<div>
+
+<p>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.</p>
+
+<p>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.</p>
+
+<p>Code Generator Options:</p>
+<table border="1" cellspacing="0">
+ <tr>
+ <th>Option</th>
+ <th>Description</th>
+ </tr>
+ <tr>
+ <td><code>double</code></td>
+ <td align="left">If enabled, the map_f64_to_f32 directive is
+ disabled in the PTX output, allowing native double-precision
+ arithmetic</td>
+ </tr>
+ <tr>
+ <td><code>no-fma</code></td>
+ <td align="left">Disable generation of Fused-Multiply Add
+ instructions, which may be beneficial for some devices</td>
+ </tr>
+ <tr>
+ <td><code>smxy / computexy</code></td>
+ <td align="left">Set shader model/compute capability to x.y,
+ e.g. sm20 or compute13</td>
+ </tr>
+</table>
+
+<p>Working:</p>
+<ul>
+ <li>Arithmetic instruction selection (including combo FMA)</li>
+ <li>Bitwise instruction selection</li>
+ <li>Control-flow instruction selection</li>
+ <li>Function calls (only on SM 2.0+ and no return arguments)</li>
+ <li>Addresses spaces (0 = global, 1 = constant, 2 = local, 4 =
+ shared)</li>
+ <li>Thread synchronization (bar.sync)</li>
+ <li>Special register reads ([N]TID, [N]CTAID, PMx, CLOCK, etc.)</li>
+</ul>
+
+<p>In Progress:</p>
+<ul>
+ <li>Robust call instruction selection</li>
+ <li>Stack frame allocation</li>
+ <li>Device-specific instruction scheduling optimizations</li>
+</ul>
+
+
</div>
</div>
projects / oota-llvm.git / blobdiff