Add a natural stack alignment field to TargetData, and prevent InstCombine from

[oota-llvm.git] / docs / CodeGenerator.html
diff --git a/docs/CodeGenerator.html b/docs/CodeGenerator.html

index 4d0813fd6dc82dcb41430beb1435eaadbf4a3f6b..ca2bb6074610782e9dacd5b2e90b308f56bd5d7f 100644 (file)
--- a/docs/CodeGenerator.html
+++ b/docs/CodeGenerator.html
@@ -5,12 +5,23 @@
    <meta http-equiv="content-type" content="text/html; charset=utf-8">
    <title>The LLVM Target-Independent Code Generator</title>
    <link rel="stylesheet" href="llvm.css" type="text/css">
    <meta http-equiv="content-type" content="text/html; charset=utf-8">
    <title>The LLVM Target-Independent Code Generator</title>
    <link rel="stylesheet" href="llvm.css" type="text/css">
+
+  <style type="text/css">
+    .unknown { background-color: #C0C0C0; text-align: center; }
+    .unknown:before { content: "?" }
+    .no { background-color: #C11B17 }
+    .no:before { content: "N" }
+    .partial { background-color: #F88017 }
+    .yes { background-color: #0F0; }
+    .yes:before { content: "Y" }
+  </style>
+
  </head>
  <body>
  
  </head>
  <body>
  
-<div class="doc_title">
+<h1>
    The LLVM Target-Independent Code Generator
    The LLVM Target-Independent Code Generator
-</div>
+</h1>
  
  <ol>
    <li><a href="#introduction">Introduction</a>
  
  <ol>
    <li><a href="#introduction">Introduction</a>
@@ -33,7 +44,7 @@
        <li><a href="#targetjitinfo">The <tt>TargetJITInfo</tt> class</a></li>
      </ul>
    </li>
        <li><a href="#targetjitinfo">The <tt>TargetJITInfo</tt> class</a></li>
      </ul>
    </li>
-  <li><a href="#codegendesc">Machine code description classes</a>
+  <li><a href="#codegendesc">The "Machine" Code Generator classes</a>
      <ul>
      <li><a href="#machineinstr">The <tt>MachineInstr</tt> class</a></li>
      <li><a href="#machinebasicblock">The <tt>MachineBasicBlock</tt>
      <ul>
      <li><a href="#machineinstr">The <tt>MachineInstr</tt> class</a></li>
      <li><a href="#machinebasicblock">The <tt>MachineBasicBlock</tt>
@@ -41,6 +52,15 @@
      <li><a href="#machinefunction">The <tt>MachineFunction</tt> class</a></li>
      </ul>
    </li>
      <li><a href="#machinefunction">The <tt>MachineFunction</tt> class</a></li>
      </ul>
    </li>
+  <li><a href="#mc">The "MC" Layer</a>
+    <ul>
+    <li><a href="#mcstreamer">The <tt>MCStreamer</tt> API</a></li>
+    <li><a href="#mccontext">The <tt>MCContext</tt> class</a>
+    <li><a href="#mcsymbol">The <tt>MCSymbol</tt> class</a></li>
+    <li><a href="#mcsection">The <tt>MCSection</tt> class</a></li>
+    <li><a href="#mcinst">The <tt>MCInst</tt> class</a></li>
+    </ul>
+  </li>
    <li><a href="#codegenalgs">Target-independent code generation algorithms</a>
      <ul>
      <li><a href="#instselect">Instruction Selection</a>
    <li><a href="#codegenalgs">Target-independent code generation algorithms</a>
      <ul>
      <li><a href="#instselect">Instruction Selection</a>
@@ -76,15 +96,14 @@
        <li><a href="#regAlloc_fold">Instruction folding</a></li>
        <li><a href="#regAlloc_builtIn">Built in register allocators</a></li>
        </ul></li>
        <li><a href="#regAlloc_fold">Instruction folding</a></li>
        <li><a href="#regAlloc_builtIn">Built in register allocators</a></li>
        </ul></li>
-    <li><a href="#codeemit">Code Emission</a>
-        <ul>
-        <li><a href="#codeemit_asm">Generating Assembly Code</a></li>
-        <li><a href="#codeemit_bin">Generating Binary Machine Code</a></li>
-        </ul></li>
+    <li><a href="#codeemit">Code Emission</a></li>
      </ul>
    </li>
      </ul>
    </li>
+  <li><a href="#nativeassembler">Implementing a Native Assembler</a></li>
+  
    <li><a href="#targetimpls">Target-specific Implementation Notes</a>
      <ul>
    <li><a href="#targetimpls">Target-specific Implementation Notes</a>
      <ul>
+    <li><a href="#targetfeatures">Target Feature Matrix</a></li>
      <li><a href="#tailcallopt">Tail call optimization</a></li>
      <li><a href="#sibcallopt">Sibling call optimization</a></li>
      <li><a href="#x86">The X86 backend</a></li>
      <li><a href="#tailcallopt">Tail call optimization</a></li>
      <li><a href="#sibcallopt">Sibling call optimization</a></li>
      <li><a href="#x86">The X86 backend</a></li>
@@ -95,16 +114,13 @@
        <li><a href="#ppc_prolog">Prolog/Epilog</a></li>
        <li><a href="#ppc_dynamic">Dynamic Allocation</a></li>
        </ul></li>
        <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>
  
  <div class="doc_author">
      </ul></li>
  
  </ol>
  
  <div class="doc_author">
-  <p>Written by <a href="mailto:sabre@nondot.org">Chris Lattner</a>,
-                <a href="mailto:isanbard@gmail.com">Bill Wendling</a>,
-                <a href="mailto:pronesto@gmail.com">Fernando Magno Quintao
-                                                    Pereira</a> and
-                <a href="mailto:jlaskey@mac.com">Jim Laskey</a></p>
+  <p>Written by the LLVM Team.</p>
  </div>
  
  <div class="doc_warning">
  </div>
  
  <div class="doc_warning">
@@ -112,18 +128,18 @@
  </div>
  
  <!-- *********************************************************************** -->
  </div>
  
  <!-- *********************************************************************** -->
-<div class="doc_section">
+<h2>
    <a name="introduction">Introduction</a>
    <a name="introduction">Introduction</a>
-</div>
+</h2>
  <!-- *********************************************************************** -->
  
  <!-- *********************************************************************** -->
  
-<div class="doc_text">
+<div>
  
  <p>The LLVM target-independent code generator is a framework that provides a
     suite of reusable components for translating the LLVM internal representation
     to the machine code for a specified target&mdash;either in assembly form
     (suitable for a static compiler) or in binary machine code format (usable for
  
  <p>The LLVM target-independent code generator is a framework that provides a
     suite of reusable components for translating the LLVM internal representation
     to the machine code for a specified target&mdash;either in assembly form
     (suitable for a static compiler) or in binary machine code format (usable for
-   a JIT compiler). The LLVM target-independent code generator consists of five
+   a JIT compiler). The LLVM target-independent code generator consists of six
     main components:</p>
  
  <ol>
     main components:</p>
  
  <ol>
@@ -132,10 +148,17 @@
        independently of how they will be used.  These interfaces are defined in
        <tt>include/llvm/Target/</tt>.</li>
  
        independently of how they will be used.  These interfaces are defined in
        <tt>include/llvm/Target/</tt>.</li>
  
-  <li>Classes used to represent the <a href="#codegendesc">machine code</a>
-      being generated for a target.  These classes are intended to be abstract
+  <li>Classes used to represent the <a href="#codegendesc">code being
+      generated</a> for a target.  These classes are intended to be abstract
        enough to represent the machine code for <i>any</i> target machine.  These
        enough to represent the machine code for <i>any</i> target machine.  These
-      classes are defined in <tt>include/llvm/CodeGen/</tt>.</li>
+      classes are defined in <tt>include/llvm/CodeGen/</tt>. At this level,
+      concepts like "constant pool entries" and "jump tables" are explicitly
+      exposed.</li>
+
+  <li>Classes and algorithms used to represent code as the object file level,
+      the <a href="#mc">MC Layer</a>.  These classes represent assembly level
+      constructs like labels, sections, and instructions.  At this level,
+      concepts like "constant pool entries" and "jump tables" don't exist.</li>
  
    <li><a href="#codegenalgs">Target-independent algorithms</a> used to implement
        various phases of native code generation (register allocation, scheduling,
  
    <li><a href="#codegenalgs">Target-independent algorithms</a> used to implement
        various phases of native code generation (register allocation, scheduling,
@@ -166,14 +189,12 @@
     depend on the target-description and machine code representation classes,
     ensuring that it is portable.</p>
  
     depend on the target-description and machine code representation classes,
     ensuring that it is portable.</p>
  
-</div>
-
  <!-- ======================================================================= -->
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
   <a name="required">Required components in the code generator</a>
   <a name="required">Required components in the code generator</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The two pieces of the LLVM code generator are the high-level interface to the
     code generator and the set of reusable components that can be used to build
  
  <p>The two pieces of the LLVM code generator are the high-level interface to the
     code generator and the set of reusable components that can be used to build
@@ -201,11 +222,11 @@
  </div>
  
  <!-- ======================================================================= -->
  </div>
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
   <a name="high-level-design">The high-level design of the code generator</a>
   <a name="high-level-design">The high-level design of the code generator</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The LLVM target-independent code generator is designed to support efficient
     and quality code generation for standard register-based microprocessors.
  
  <p>The LLVM target-independent code generator is designed to support efficient
     and quality code generation for standard register-based microprocessors.
@@ -275,11 +296,11 @@
  </div>
  
  <!-- ======================================================================= -->
  </div>
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
   <a name="tablegen">Using TableGen for target description</a>
   <a name="tablegen">Using TableGen for target description</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The target description classes require a detailed description of the target
     architecture.  These target descriptions often have a large amount of common
  
  <p>The target description classes require a detailed description of the target
     architecture.  These target descriptions often have a large amount of common
@@ -302,13 +323,15 @@
  
  </div>
  
  
  </div>
  
+</div>
+
  <!-- *********************************************************************** -->
  <!-- *********************************************************************** -->
-<div class="doc_section">
+<h2>
    <a name="targetdesc">Target description classes</a>
    <a name="targetdesc">Target description classes</a>
-</div>
+</h2>
  <!-- *********************************************************************** -->
  
  <!-- *********************************************************************** -->
  
-<div class="doc_text">
+<div>
  
  <p>The LLVM target description classes (located in the
     <tt>include/llvm/Target</tt> directory) provide an abstract description of
  
  <p>The LLVM target description classes (located in the
     <tt>include/llvm/Target</tt> directory) provide an abstract description of
@@ -324,14 +347,12 @@
     <tt><a href="#targetmachine">TargetMachine</a></tt> class provides accessors
     that should be implemented by the target.</p>
  
     <tt><a href="#targetmachine">TargetMachine</a></tt> class provides accessors
     that should be implemented by the target.</p>
  
-</div>
-
  <!-- ======================================================================= -->
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="targetmachine">The <tt>TargetMachine</tt> class</a>
    <a name="targetmachine">The <tt>TargetMachine</tt> class</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The <tt>TargetMachine</tt> class provides virtual methods that are used to
     access the target-specific implementations of the various target description
  
  <p>The <tt>TargetMachine</tt> class provides virtual methods that are used to
     access the target-specific implementations of the various target description
@@ -347,11 +368,11 @@
  </div>
  
  <!-- ======================================================================= -->
  </div>
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="targetdata">The <tt>TargetData</tt> class</a>
    <a name="targetdata">The <tt>TargetData</tt> class</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The <tt>TargetData</tt> class is the only required target description class,
     and it is the only class that is not extensible (you cannot derived a new
  
  <p>The <tt>TargetData</tt> class is the only required target description class,
     and it is the only class that is not extensible (you cannot derived a new
@@ -363,11 +384,11 @@
  </div>
  
  <!-- ======================================================================= -->
  </div>
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="targetlowering">The <tt>TargetLowering</tt> class</a>
    <a name="targetlowering">The <tt>TargetLowering</tt> class</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The <tt>TargetLowering</tt> class is used by SelectionDAG based instruction
     selectors primarily to describe how LLVM code should be lowered to
  
  <p>The <tt>TargetLowering</tt> class is used by SelectionDAG based instruction
     selectors primarily to describe how LLVM code should be lowered to
@@ -389,11 +410,11 @@
  </div>
  
  <!-- ======================================================================= -->
  </div>
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="targetregisterinfo">The <tt>TargetRegisterInfo</tt> class</a>
    <a name="targetregisterinfo">The <tt>TargetRegisterInfo</tt> class</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The <tt>TargetRegisterInfo</tt> class is used to describe the register file
     of the target and any interactions between the registers.</p>
  
  <p>The <tt>TargetRegisterInfo</tt> class is used to describe the register file
     of the target and any interactions between the registers.</p>
@@ -423,11 +444,11 @@
  </div>
  
  <!-- ======================================================================= -->
  </div>
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="targetinstrinfo">The <tt>TargetInstrInfo</tt> class</a>
    <a name="targetinstrinfo">The <tt>TargetInstrInfo</tt> class</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The <tt>TargetInstrInfo</tt> class is used to describe the machine
     instructions supported by the target. It is essentially an array of
  
  <p>The <tt>TargetInstrInfo</tt> class is used to describe the machine
     instructions supported by the target. It is essentially an array of
@@ -441,11 +462,11 @@
  </div>
  
  <!-- ======================================================================= -->
  </div>
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="targetframeinfo">The <tt>TargetFrameInfo</tt> class</a>
    <a name="targetframeinfo">The <tt>TargetFrameInfo</tt> class</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The <tt>TargetFrameInfo</tt> class is used to provide information about the
     stack frame layout of the target. It holds the direction of stack growth, the
  
  <p>The <tt>TargetFrameInfo</tt> class is used to provide information about the
     stack frame layout of the target. It holds the direction of stack growth, the
@@ -457,11 +478,11 @@
  </div>
  
  <!-- ======================================================================= -->
  </div>
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="targetsubtarget">The <tt>TargetSubtarget</tt> class</a>
    <a name="targetsubtarget">The <tt>TargetSubtarget</tt> class</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The <tt>TargetSubtarget</tt> class is used to provide information about the
     specific chip set being targeted.  A sub-target informs code generation of
  
  <p>The <tt>TargetSubtarget</tt> class is used to provide information about the
     specific chip set being targeted.  A sub-target informs code generation of
@@ -473,11 +494,11 @@
  
  
  <!-- ======================================================================= -->
  
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="targetjitinfo">The <tt>TargetJITInfo</tt> class</a>
    <a name="targetjitinfo">The <tt>TargetJITInfo</tt> class</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The <tt>TargetJITInfo</tt> class exposes an abstract interface used by the
     Just-In-Time code generator to perform target-specific activities, such as
  
  <p>The <tt>TargetJITInfo</tt> class exposes an abstract interface used by the
     Just-In-Time code generator to perform target-specific activities, such as
@@ -487,13 +508,15 @@
  
  </div>
  
  
  </div>
  
+</div>
+
  <!-- *********************************************************************** -->
  <!-- *********************************************************************** -->
-<div class="doc_section">
+<h2>
    <a name="codegendesc">Machine code description classes</a>
    <a name="codegendesc">Machine code description classes</a>
-</div>
+</h2>
  <!-- *********************************************************************** -->
  
  <!-- *********************************************************************** -->
  
-<div class="doc_text">
+<div>
  
  <p>At the high-level, LLVM code is translated to a machine specific
     representation formed out of
  
  <p>At the high-level, LLVM code is translated to a machine specific
     representation formed out of
@@ -506,14 +529,12 @@
     SSA representation for machine code, as well as a register allocated, non-SSA
     form.</p>
  
     SSA representation for machine code, as well as a register allocated, non-SSA
     form.</p>
  
-</div>
-
  <!-- ======================================================================= -->
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="machineinstr">The <tt>MachineInstr</tt> class</a>
    <a name="machineinstr">The <tt>MachineInstr</tt> class</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>Target machine instructions are represented as instances of the
     <tt>MachineInstr</tt> class.  This class is an extremely abstract way of
  
  <p>Target machine instructions are represented as instances of the
     <tt>MachineInstr</tt> class.  This class is an extremely abstract way of
@@ -554,14 +575,12 @@
  <p>Also if the first operand is a def, it is easier to <a href="#buildmi">create
     instructions</a> whose only def is the first operand.</p>
  
  <p>Also if the first operand is a def, it is easier to <a href="#buildmi">create
     instructions</a> whose only def is the first operand.</p>
  
-</div>
-
  <!-- _______________________________________________________________________ -->
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="buildmi">Using the <tt>MachineInstrBuilder.h</tt> functions</a>
    <a name="buildmi">Using the <tt>MachineInstrBuilder.h</tt> functions</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>Machine instructions are created by using the <tt>BuildMI</tt> functions,
     located in the <tt>include/llvm/CodeGen/MachineInstrBuilder.h</tt> file.  The
  
  <p>Machine instructions are created by using the <tt>BuildMI</tt> functions,
     located in the <tt>include/llvm/CodeGen/MachineInstrBuilder.h</tt> file.  The
@@ -608,11 +627,11 @@ MI.addReg(Reg, RegState::Define);
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="fixedregs">Fixed (preassigned) registers</a>
    <a name="fixedregs">Fixed (preassigned) registers</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>One important issue that the code generator needs to be aware of is the
     presence of fixed registers.  In particular, there are often places in the
  
  <p>One important issue that the code generator needs to be aware of is the
     presence of fixed registers.  In particular, there are often places in the
@@ -680,11 +699,11 @@ ret
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="ssa">Machine code in SSA form</a>
    <a name="ssa">Machine code in SSA form</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p><tt>MachineInstr</tt>'s are initially selected in SSA-form, and are
     maintained in SSA-form until register allocation happens.  For the most part,
  
  <p><tt>MachineInstr</tt>'s are initially selected in SSA-form, and are
     maintained in SSA-form until register allocation happens.  For the most part,
@@ -697,12 +716,14 @@ ret
  
  </div>
  
  
  </div>
  
+</div>
+
  <!-- ======================================================================= -->
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="machinebasicblock">The <tt>MachineBasicBlock</tt> class</a>
    <a name="machinebasicblock">The <tt>MachineBasicBlock</tt> class</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The <tt>MachineBasicBlock</tt> class contains a list of machine instructions
     (<tt><a href="#machineinstr">MachineInstr</a></tt> instances).  It roughly
  
  <p>The <tt>MachineBasicBlock</tt> class contains a list of machine instructions
     (<tt><a href="#machineinstr">MachineInstr</a></tt> instances).  It roughly
@@ -715,11 +736,11 @@ ret
  </div>
  
  <!-- ======================================================================= -->
  </div>
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="machinefunction">The <tt>MachineFunction</tt> class</a>
    <a name="machinefunction">The <tt>MachineFunction</tt> class</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The <tt>MachineFunction</tt> class contains a list of machine basic blocks
     (<tt><a href="#machinebasicblock">MachineBasicBlock</a></tt> instances).  It
  
  <p>The <tt>MachineFunction</tt> class contains a list of machine basic blocks
     (<tt><a href="#machinebasicblock">MachineBasicBlock</a></tt> instances).  It
@@ -732,26 +753,174 @@ ret
  
  </div>
  
  
  </div>
  
+</div>
+
  <!-- *********************************************************************** -->
  <!-- *********************************************************************** -->
-<div class="doc_section">
-  <a name="codegenalgs">Target-independent code generation algorithms</a>
+<h2>
+  <a name="mc">The "MC" Layer</a>
+</h2>
+<!-- *********************************************************************** -->
+
+<div>
+
+<p>
+The MC Layer is used to represent and process code at the raw machine code
+level, devoid of "high level" information like "constant pools", "jump tables",
+"global variables" or anything like that.  At this level, LLVM handles things
+like label names, machine instructions, and sections in the object file.  The
+code in this layer is used for a number of important purposes: the tail end of
+the code generator uses it to write a .s or .o file, and it is also used by the
+llvm-mc tool to implement standalone machine code assemblers and disassemblers.
+</p>
+
+<p>
+This section describes some of the important classes.  There are also a number
+of important subsystems that interact at this layer, they are described later
+in this manual.
+</p>
+
+<!-- ======================================================================= -->
+<h3>
+  <a name="mcstreamer">The <tt>MCStreamer</tt> API</a>
+</h3>
+
+<div>
+
+<p>
+MCStreamer is best thought of as an assembler API.  It is an abstract API which
+is <em>implemented</em> in different ways (e.g. to output a .s file, output an
+ELF .o file, etc) but whose API correspond directly to what you see in a .s
+file.  MCStreamer has one method per directive, such as EmitLabel,
+EmitSymbolAttribute, SwitchSection, EmitValue (for .byte, .word), etc, which
+directly correspond to assembly level directives.  It also has an
+EmitInstruction method, which is used to output an MCInst to the streamer.
+</p>
+
+<p>
+This API is most important for two clients: the llvm-mc stand-alone assembler is
+effectively a parser that parses a line, then invokes a method on MCStreamer. In
+the code generator, the <a href="#codeemit">Code Emission</a> phase of the code
+generator lowers higher level LLVM IR and Machine* constructs down to the MC
+layer, emitting directives through MCStreamer.</p>
+
+<p>
+On the implementation side of MCStreamer, there are two major implementations:
+one for writing out a .s file (MCAsmStreamer), and one for writing out a .o
+file (MCObjectStreamer).  MCAsmStreamer is a straight-forward implementation
+that prints out a directive for each method (e.g. EmitValue -&gt; .byte), but
+MCObjectStreamer implements a full assembler.
+</p>
+
+</div>
+
+<!-- ======================================================================= -->
+<h3>
+  <a name="mccontext">The <tt>MCContext</tt> class</a>
+</h3>
+
+<div>
+
+<p>
+The MCContext class is the owner of a variety of uniqued data structures at the
+MC layer, including symbols, sections, etc.  As such, this is the class that you
+interact with to create symbols and sections.  This class can not be subclassed.
+</p>
+
+</div>
+
+<!-- ======================================================================= -->
+<h3>
+  <a name="mcsymbol">The <tt>MCSymbol</tt> class</a>
+</h3>
+
+<div>
+
+<p>
+The MCSymbol class represents a symbol (aka label) in the assembly file.  There
+are two interesting kinds of symbols: assembler temporary symbols, and normal
+symbols.  Assembler temporary symbols are used and processed by the assembler
+but are discarded when the object file is produced.  The distinction is usually
+represented by adding a prefix to the label, for example "L" labels are
+assembler temporary labels in MachO.
+</p>
+
+<p>MCSymbols are created by MCContext and uniqued there.  This means that
+MCSymbols can be compared for pointer equivalence to find out if they are the
+same symbol.  Note that pointer inequality does not guarantee the labels will
+end up at different addresses though.  It's perfectly legal to output something
+like this to the .s file:<p>
+
+<pre>
+  foo:
+  bar:
+    .byte 4
+</pre>
+
+<p>In this case, both the foo and bar symbols will have the same address.</p>
+
+</div>
+
+<!-- ======================================================================= -->
+<h3>
+  <a name="mcsection">The <tt>MCSection</tt> class</a>
+</h3>
+
+<div>
+
+<p>
+The MCSection class represents an object-file specific section. It is subclassed
+by object file specific implementations (e.g. <tt>MCSectionMachO</tt>, 
+<tt>MCSectionCOFF</tt>, <tt>MCSectionELF</tt>) and these are created and uniqued
+by MCContext.  The MCStreamer has a notion of the current section, which can be
+changed with the SwitchToSection method (which corresponds to a ".section"
+directive in a .s file).
+</p>
+
+</div>
+
+<!-- ======================================================================= -->
+<h3>
+  <a name="mcinst">The <tt>MCInst</tt> class</a>
+</h3>
+
+<div>
+
+<p>
+The MCInst class is a target-independent representation of an instruction.  It
+is a simple class (much more so than <a href="#machineinstr">MachineInstr</a>)
+that holds a target-specific opcode and a vector of MCOperands.  MCOperand, in
+turn, is a simple discriminated union of three cases: 1) a simple immediate, 
+2) a target register ID, 3) a symbolic expression (e.g. "Lfoo-Lbar+42") as an
+MCExpr.
+</p>
+
+<p>MCInst is the common currency used to represent machine instructions at the
+MC layer.  It is the type used by the instruction encoder, the instruction
+printer, and the type generated by the assembly parser and disassembler.
+</p>
+
+</div>
+
  </div>
  </div>
+
+<!-- *********************************************************************** -->
+<h2>
+  <a name="codegenalgs">Target-independent code generation algorithms</a>
+</h2>
  <!-- *********************************************************************** -->
  
  <!-- *********************************************************************** -->
  
-<div class="doc_text">
+<div>
  
  <p>This section documents the phases described in the
     <a href="#high-level-design">high-level design of the code generator</a>.
     It explains how they work and some of the rationale behind their design.</p>
  
  
  <p>This section documents the phases described in the
     <a href="#high-level-design">high-level design of the code generator</a>.
     It explains how they work and some of the rationale behind their design.</p>
  
-</div>
-
  <!-- ======================================================================= -->
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="instselect">Instruction Selection</a>
    <a name="instselect">Instruction Selection</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>Instruction Selection is the process of translating LLVM code presented to
     the code generator into target-specific machine instructions.  There are
  
  <p>Instruction Selection is the process of translating LLVM code presented to
     the code generator into target-specific machine instructions.  There are
@@ -763,14 +932,12 @@ ret
     selector to be generated from these <tt>.td</tt> files, though currently
     there are still things that require custom C++ code.</p>
  
     selector to be generated from these <tt>.td</tt> files, though currently
     there are still things that require custom C++ code.</p>
  
-</div>
-
  <!-- _______________________________________________________________________ -->
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="selectiondag_intro">Introduction to SelectionDAGs</a>
    <a name="selectiondag_intro">Introduction to SelectionDAGs</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The SelectionDAG provides an abstraction for code representation in a way
     that is amenable to instruction selection using automatic techniques
  
  <p>The SelectionDAG provides an abstraction for code representation in a way
     that is amenable to instruction selection using automatic techniques
@@ -828,11 +995,11 @@ ret
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="selectiondag_process">SelectionDAG Instruction Selection Process</a>
    <a name="selectiondag_process">SelectionDAG Instruction Selection Process</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>SelectionDAG-based instruction selection consists of the following steps:</p>
  
  
  <p>SelectionDAG-based instruction selection consists of the following steps:</p>
  
@@ -857,9 +1024,9 @@ ret
        SelectionDAG optimizer is run to clean up redundancies exposed by type
        legalization.</li>
  
        SelectionDAG optimizer is run to clean up redundancies exposed by type
        legalization.</li>
  
-  <li><a href="#selectiondag_legalize">Legalize SelectionDAG Types</a> &mdash;
-      This stage transforms SelectionDAG nodes to eliminate any types that are
-      unsupported on the target.</li>
+  <li><a href="#selectiondag_legalize">Legalize SelectionDAG Ops</a> &mdash;
+      This stage transforms SelectionDAG nodes to eliminate any operations 
+      that are unsupported on the target.</li>
  
    <li><a href="#selectiondag_optimize">Optimize SelectionDAG</a> &mdash; The
        SelectionDAG optimizer is run to eliminate inefficiencies introduced by
  
    <li><a href="#selectiondag_optimize">Optimize SelectionDAG</a> &mdash; The
        SelectionDAG optimizer is run to eliminate inefficiencies introduced by
@@ -909,11 +1076,11 @@ ret
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="selectiondag_build">Initial SelectionDAG Construction</a>
    <a name="selectiondag_build">Initial SelectionDAG Construction</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The initial SelectionDAG is na&iuml;vely peephole expanded from the LLVM
     input by the <tt>SelectionDAGLowering</tt> class in the
  
  <p>The initial SelectionDAG is na&iuml;vely peephole expanded from the LLVM
     input by the <tt>SelectionDAGLowering</tt> class in the
@@ -929,11 +1096,11 @@ ret
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="selectiondag_legalize_types">SelectionDAG LegalizeTypes Phase</a>
    <a name="selectiondag_legalize_types">SelectionDAG LegalizeTypes Phase</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The Legalize phase is in charge of converting a DAG to only use the types
     that are natively supported by the target.</p>
  
  <p>The Legalize phase is in charge of converting a DAG to only use the types
     that are natively supported by the target.</p>
@@ -962,11 +1129,11 @@ ret
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="selectiondag_legalize">SelectionDAG Legalize Phase</a>
    <a name="selectiondag_legalize">SelectionDAG Legalize Phase</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The Legalize phase is in charge of converting a DAG to only use the
     operations that are natively supported by the target.</p>
  
  <p>The Legalize phase is in charge of converting a DAG to only use the
     operations that are natively supported by the target.</p>
@@ -994,12 +1161,13 @@ ret
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
-  <a name="selectiondag_optimize">SelectionDAG Optimization Phase: the DAG
-  Combiner</a>
-</div>
+<h4>
+  <a name="selectiondag_optimize">
+    SelectionDAG Optimization Phase: the DAG Combiner
+  </a>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The SelectionDAG optimization phase is run multiple times for code
     generation, immediately after the DAG is built and once after each
  
  <p>The SelectionDAG optimization phase is run multiple times for code
     generation, immediately after the DAG is built and once after each
@@ -1029,11 +1197,11 @@ ret
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="selectiondag_select">SelectionDAG Select Phase</a>
    <a name="selectiondag_select">SelectionDAG Select Phase</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The Select phase is the bulk of the target-specific code for instruction
     selection.  This phase takes a legal SelectionDAG as input, pattern matches
  
  <p>The Select phase is the bulk of the target-specific code for instruction
     selection.  This phase takes a legal SelectionDAG as input, pattern matches
@@ -1190,11 +1358,11 @@ def : Pat&lt;(i32 imm:$imm),
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="selectiondag_sched">SelectionDAG Scheduling and Formation Phase</a>
    <a name="selectiondag_sched">SelectionDAG Scheduling and Formation Phase</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The scheduling phase takes the DAG of target instructions from the selection
     phase and assigns an order.  The scheduler can pick an order depending on
  
  <p>The scheduling phase takes the DAG of target instructions from the selection
     phase and assigns an order.  The scheduler can pick an order depending on
@@ -1211,11 +1379,11 @@ def : Pat&lt;(i32 imm:$imm),
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="selectiondag_future">Future directions for the SelectionDAG</a>
    <a name="selectiondag_future">Future directions for the SelectionDAG</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <ol>
    <li>Optional function-at-a-time selection.</li>
  
  <ol>
    <li>Optional function-at-a-time selection.</li>
@@ -1225,18 +1393,20 @@ def : Pat&lt;(i32 imm:$imm),
  
  </div>
   
  
  </div>
   
+</div>
+
  <!-- ======================================================================= -->
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="ssamco">SSA-based Machine Code Optimizations</a>
    <a name="ssamco">SSA-based Machine Code Optimizations</a>
-</div>
-<div class="doc_text"><p>To Be Written</p></div>
+</h3>
+<div><p>To Be Written</p></div>
  
  <!-- ======================================================================= -->
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="liveintervals">Live Intervals</a>
    <a name="liveintervals">Live Intervals</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>Live Intervals are the ranges (intervals) where a variable is <i>live</i>.
     They are used by some <a href="#regalloc">register allocator</a> passes to
  
  <p>Live Intervals are the ranges (intervals) where a variable is <i>live</i>.
     They are used by some <a href="#regalloc">register allocator</a> passes to
@@ -1244,14 +1414,12 @@ def : Pat&lt;(i32 imm:$imm),
     register are live at the same point in the program (i.e., they conflict).
     When this situation occurs, one virtual register must be <i>spilled</i>.</p>
  
     register are live at the same point in the program (i.e., they conflict).
     When this situation occurs, one virtual register must be <i>spilled</i>.</p>
  
-</div>
-
  <!-- _______________________________________________________________________ -->
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="livevariable_analysis">Live Variable Analysis</a>
    <a name="livevariable_analysis">Live Variable Analysis</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The first step in determining the live intervals of variables is to calculate
     the set of registers that are immediately dead after the instruction (i.e.,
  
  <p>The first step in determining the live intervals of variables is to calculate
     the set of registers that are immediately dead after the instruction (i.e.,
@@ -1293,11 +1461,11 @@ def : Pat&lt;(i32 imm:$imm),
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="liveintervals_analysis">Live Intervals Analysis</a>
    <a name="liveintervals_analysis">Live Intervals Analysis</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>We now have the information available to perform the live intervals analysis
     and build the live intervals themselves.  We start off by numbering the basic
  
  <p>We now have the information available to perform the live intervals analysis
     and build the live intervals themselves.  We start off by numbering the basic
@@ -1312,12 +1480,14 @@ def : Pat&lt;(i32 imm:$imm),
  
  </div>
  
  
  </div>
  
+</div>
+
  <!-- ======================================================================= -->
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="regalloc">Register Allocation</a>
    <a name="regalloc">Register Allocation</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The <i>Register Allocation problem</i> consists in mapping a program
     <i>P<sub>v</sub></i>, that can use an unbounded number of virtual registers,
  
  <p>The <i>Register Allocation problem</i> consists in mapping a program
     <i>P<sub>v</sub></i>, that can use an unbounded number of virtual registers,
@@ -1327,15 +1497,13 @@ def : Pat&lt;(i32 imm:$imm),
     accommodate all the virtual registers, some of them will have to be mapped
     into memory. These virtuals are called <i>spilled virtuals</i>.</p>
  
     accommodate all the virtual registers, some of them will have to be mapped
     into memory. These virtuals are called <i>spilled virtuals</i>.</p>
  
-</div>
-
  <!-- _______________________________________________________________________ -->
  
  <!-- _______________________________________________________________________ -->
  
-<div class="doc_subsubsection">
+<h4>
    <a name="regAlloc_represent">How registers are represented in LLVM</a>
    <a name="regAlloc_represent">How registers are represented in LLVM</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>In LLVM, physical registers are denoted by integer numbers that normally
     range from 1 to 1023. To see how this numbering is defined for a particular
  
  <p>In LLVM, physical registers are denoted by integer numbers that normally
     range from 1 to 1023. To see how this numbering is defined for a particular
@@ -1386,18 +1554,25 @@ bool RegMapping_Fer::compatible_class(MachineFunction &amp;mf,
     </p>
  
  <p>Virtual registers are also denoted by integer numbers. Contrary to physical
     </p>
  
  <p>Virtual registers are also denoted by integer numbers. Contrary to physical
-   registers, different virtual registers never share the same number. The
-   smallest virtual register is normally assigned the number 1024. This may
-   change, so, in order to know which is the first virtual register, you should
-   access <tt>TargetRegisterInfo::FirstVirtualRegister</tt>. Any register whose
-   number is greater than or equal
-   to <tt>TargetRegisterInfo::FirstVirtualRegister</tt> is considered a virtual
-   register. Whereas physical registers are statically defined in
-   a <tt>TargetRegisterInfo.td</tt> file and cannot be created by the
-   application developer, that is not the case with virtual registers.  In order
-   to create new virtual registers, use the
+   registers, different virtual registers never share the same number. Whereas
+   physical registers are statically defined in a <tt>TargetRegisterInfo.td</tt>
+   file and cannot be created by the application developer, that is not the case
+   with virtual registers. In order to create new virtual registers, use the
     method <tt>MachineRegisterInfo::createVirtualRegister()</tt>. This method
     method <tt>MachineRegisterInfo::createVirtualRegister()</tt>. This method
-   will return a virtual register with the highest code.</p>
+   will return a new virtual register. Use an <tt>IndexedMap&lt;Foo,
+   VirtReg2IndexFunctor&gt;</tt> to hold information per virtual register. If you
+   need to enumerate all virtual registers, use the function
+   <tt>TargetRegisterInfo::index2VirtReg()</tt> to find the virtual register
+   numbers:</p>
+
+<div class="doc_code">
+<pre>
+  for (unsigned i = 0, e = MRI->getNumVirtRegs(); i != e; ++i) {
+    unsigned VirtReg = TargetRegisterInfo::index2VirtReg(i);
+    stuff(VirtReg);
+  }
+</pre>
+</div>
  
  <p>Before register allocation, the operands of an instruction are mostly virtual
     registers, although physical registers may also be used. In order to check if
  
  <p>Before register allocation, the operands of an instruction are mostly virtual
     registers, although physical registers may also be used. In order to check if
@@ -1437,11 +1612,11 @@ bool RegMapping_Fer::compatible_class(MachineFunction &amp;mf,
  
  <!-- _______________________________________________________________________ -->
  
  
  <!-- _______________________________________________________________________ -->
  
-<div class="doc_subsubsection">
+<h4>
    <a name="regAlloc_howTo">Mapping virtual registers to physical registers</a>
    <a name="regAlloc_howTo">Mapping virtual registers to physical registers</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>There are two ways to map virtual registers to physical registers (or to
     memory slots). The first way, that we will call <i>direct mapping</i>, is
  
  <p>There are two ways to map virtual registers to physical registers (or to
     memory slots). The first way, that we will call <i>direct mapping</i>, is
@@ -1487,11 +1662,11 @@ bool RegMapping_Fer::compatible_class(MachineFunction &amp;mf,
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="regAlloc_twoAddr">Handling two address instructions</a>
    <a name="regAlloc_twoAddr">Handling two address instructions</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>With very rare exceptions (e.g., function calls), the LLVM machine code
     instructions are three address instructions. That is, each instruction is
  
  <p>With very rare exceptions (e.g., function calls), the LLVM machine code
     instructions are three address instructions. That is, each instruction is
@@ -1523,11 +1698,11 @@ bool RegMapping_Fer::compatible_class(MachineFunction &amp;mf,
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="regAlloc_ssaDecon">The SSA deconstruction phase</a>
    <a name="regAlloc_ssaDecon">The SSA deconstruction phase</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>An important transformation that happens during register allocation is called
     the <i>SSA Deconstruction Phase</i>. The SSA form simplifies many analyses
  
  <p>An important transformation that happens during register allocation is called
     the <i>SSA Deconstruction Phase</i>. The SSA form simplifies many analyses
@@ -1547,11 +1722,11 @@ bool RegMapping_Fer::compatible_class(MachineFunction &amp;mf,
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="regAlloc_fold">Instruction folding</a>
    <a name="regAlloc_fold">Instruction folding</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p><i>Instruction folding</i> is an optimization performed during register
     allocation that removes unnecessary copy instructions. For instance, a
  
  <p><i>Instruction folding</i> is an optimization performed during register
     allocation that removes unnecessary copy instructions. For instance, a
@@ -1584,32 +1759,38 @@ bool RegMapping_Fer::compatible_class(MachineFunction &amp;mf,
  
  <!-- _______________________________________________________________________ -->
  
  
  <!-- _______________________________________________________________________ -->
  
-<div class="doc_subsubsection">
+<h4>
    <a name="regAlloc_builtIn">Built in register allocators</a>
    <a name="regAlloc_builtIn">Built in register allocators</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The LLVM infrastructure provides the application developer with three
     different register allocators:</p>
  
  <ul>
  
  <p>The LLVM infrastructure provides the application developer with three
     different register allocators:</p>
  
  <ul>
-  <li><i>Linear Scan</i> &mdash; <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> &mdash; 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>Fast</i> &mdash; 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> &mdash; 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> &mdash; <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> &mdash; 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>
    <li><i>PBQP</i> &mdash; 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
  </ul>
  
  <p>The type of register allocator used in <tt>llc</tt> can be chosen with the
@@ -1625,55 +1806,632 @@ $ llc -regalloc=pbqp file.bc -o pbqp.s;
  
  </div>
  
  
  </div>
  
+</div>
+
  <!-- ======================================================================= -->
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="proepicode">Prolog/Epilog Code Insertion</a>
    <a name="proepicode">Prolog/Epilog Code Insertion</a>
+</h3>
+
+<!-- _______________________________________________________________________ -->
+<h4>
+  <a name="compact_unwind">Compact Unwind</a>
+</h4>
+
+<div>
+
+<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 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
+   function, that compact unwind info will be encoded in
+   the <code>__TEXT,__unwind_info</code> section. If we emit DWARF unwind info,
+   the <code>__TEXT,__unwind_info</code> section will contain the offset of the
+   FDE in the <code>__TEXT,__eh_frame</code> section in the final linked
+   image.</p>
+
+<p>For X86, there are three modes for the compact unwind encoding:</p>
+
+<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,
+      then <code>ESP/RSP</code> is moved to <code>EBP/RBP</code>. Thus to
+      unwind, <code>ESP/RSP</code> is restored with the
+      current <code>EBP/RBP</code> value, then <code>EBP/RBP</code> is restored
+      by popping the stack, and the return is done by popping the stack once
+      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 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>
+    <th>i386 Register</th>
+    <th>x86-64 Regiser</th>
+  </tr>
+  <tr>
+    <td>1</td>
+    <td><code>EBX</code></td>
+    <td><code>RBX</code></td>
+  </tr>
+  <tr>
+    <td>2</td>
+    <td><code>ECX</code></td>
+    <td><code>R12</code></td>
+  </tr>
+  <tr>
+    <td>3</td>
+    <td><code>EDX</code></td>
+    <td><code>R13</code></td>
+  </tr>
+  <tr>
+    <td>4</td>
+    <td><code>EDI</code></td>
+    <td><code>R14</code></td>
+  </tr>
+  <tr>
+    <td>5</td>
+    <td><code>ESI</code></td>
+    <td><code>R15</code></td>
+  </tr>
+  <tr>
+    <td>6</td>
+    <td><code>EBP</code></td>
+    <td><code>RBP</code></td>
+  </tr>
+</table>
+
+</dd>
+
+  <dt><i>Frameless with a Small Constant Stack Size (<code>EBP</code>
+         or <code>RBP</code> is not used as a frame pointer)</i></dt>
+  <dd><p>To return, a constant (encoded in the compact unwind encoding) is added
+      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 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>
+         or <code>RBP</code> is not used as a frame pointer)</i></dt>
+  <dd><p>This case is like the "Frameless with a Small Constant Stack Size"
+      case, but the stack size is too large to encode in the compact unwind
+      encoding. Instead it requires that the function contains "<code>subl
+      $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>
+</dl>
+
  </div>
  </div>
-<div class="doc_text"><p>To Be Written</p></div>
+
  <!-- ======================================================================= -->
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="latemco">Late Machine Code Optimizations</a>
    <a name="latemco">Late Machine Code Optimizations</a>
-</div>
-<div class="doc_text"><p>To Be Written</p></div>
+</h3>
+<div><p>To Be Written</p></div>
+
  <!-- ======================================================================= -->
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="codeemit">Code Emission</a>
    <a name="codeemit">Code Emission</a>
+</h3>
+
+<div>
+
+<p>The code emission step of code generation is responsible for lowering from
+the code generator abstractions (like <a 
+href="#machinefunction">MachineFunction</a>, <a 
+href="#machineinstr">MachineInstr</a>, etc) down
+to the abstractions used by the MC layer (<a href="#mcinst">MCInst</a>, 
+<a href="#mcstreamer">MCStreamer</a>, etc).  This is
+done with a combination of several different classes: the (misnamed)
+target-independent AsmPrinter class, target-specific subclasses of AsmPrinter
+(such as SparcAsmPrinter), and the TargetLoweringObjectFile class.</p>
+
+<p>Since the MC layer works at the level of abstraction of object files, it
+doesn't have a notion of functions, global variables etc.  Instead, it thinks
+about labels, directives, and instructions.  A key class used at this time is
+the MCStreamer class.  This is an abstract API that is implemented in different
+ways (e.g. to output a .s file, output an ELF .o file, etc) that is effectively
+an "assembler API".  MCStreamer has one method per directive, such as EmitLabel,
+EmitSymbolAttribute, SwitchSection, etc, which directly correspond to assembly
+level directives.
+</p>
+
+<p>If you are interested in implementing a code generator for a target, there
+are three important things that you have to implement for your target:</p>
+
+<ol>
+<li>First, you need a subclass of AsmPrinter for your target.  This class
+implements the general lowering process converting MachineFunction's into MC
+label constructs.  The AsmPrinter base class provides a number of useful methods
+and routines, and also allows you to override the lowering process in some
+important ways.  You should get much of the lowering for free if you are
+implementing an ELF, COFF, or MachO target, because the TargetLoweringObjectFile
+class implements much of the common logic.</li>
+
+<li>Second, you need to implement an instruction printer for your target.  The
+instruction printer takes an <a href="#mcinst">MCInst</a> and renders it to a
+raw_ostream as text.  Most of this is automatically generated from the .td file
+(when you specify something like "<tt>add $dst, $src1, $src2</tt>" in the
+instructions), but you need to implement routines to print operands.</li>
+
+<li>Third, you need to implement code that lowers a <a
+href="#machineinstr">MachineInstr</a> to an MCInst, usually implemented in
+"&lt;target&gt;MCInstLower.cpp".  This lowering process is often target
+specific, and is responsible for turning jump table entries, constant pool
+indices, global variable addresses, etc into MCLabels as appropriate.  This
+translation layer is also responsible for expanding pseudo ops used by the code
+generator into the actual machine instructions they correspond to. The MCInsts
+that are generated by this are fed into the instruction printer or the encoder.
+</li>
+
+</ol>
+
+<p>Finally, at your choosing, you can also implement an subclass of
+MCCodeEmitter which lowers MCInst's into machine code bytes and relocations.
+This is important if you want to support direct .o file emission, or would like
+to implement an assembler for your target.</p>
+
  </div>
  </div>
-<div class="doc_text"><p>To Be Written</p></div>
+
+</div>
+
+<!-- *********************************************************************** -->
+<h2>
+  <a name="nativeassembler">Implementing a Native Assembler</a>
+</h2>
+<!-- *********************************************************************** -->
+
+<div>
+
+<p>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.
+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
+compiler.</p>
+
+<!-- ======================================================================= -->
+<h3 id="na_instparsing">Instruction Parsing</h3>
+
+<div><p>To Be Written</p></div>
+
+
+<!-- ======================================================================= -->
+<h3 id="na_instaliases">
+  Instruction Alias Processing
+</h3>
+
+<div>
+<p>Once the instruction is parsed, it enters the MatchInstructionImpl function.
+The MatchInstructionImpl function performs alias processing and then does
+actual matching.</p>
+
+<p>Alias processing is the phase that canonicalizes different lexical forms of
+the same instructions down to one representation.  There are several different
+kinds of alias that are possible to implement and they are listed below in the
+order that they are processed (which is in order from simplest/weakest to most
+complex/powerful).  Generally you want to use the first alias mechanism that
+meets the needs of your instruction, because it will allow a more concise
+description.</p>
+
  <!-- _______________________________________________________________________ -->
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
-  <a name="codeemit_asm">Generating Assembly Code</a>
+<h4>Mnemonic Aliases</h4>
+
+<div>
+
+<p>The first phase of alias processing is simple instruction mnemonic
+remapping for classes of instructions which are allowed with two different
+mnemonics.  This phase is a simple and unconditionally remapping from one input
+mnemonic to one output mnemonic.  It isn't possible for this form of alias to
+look at the operands at all, so the remapping must apply for all forms of a
+given mnemonic.  Mnemonic aliases are defined simply, for example X86 has:
+</p>
+
+<div class="doc_code">
+<pre>
+def : MnemonicAlias&lt;"cbw",     "cbtw"&gt;;
+def : MnemonicAlias&lt;"smovq",   "movsq"&gt;;
+def : MnemonicAlias&lt;"fldcww",  "fldcw"&gt;;
+def : MnemonicAlias&lt;"fucompi", "fucomip"&gt;;
+def : MnemonicAlias&lt;"ud2a",    "ud2"&gt;;
+</pre>
+</div>
+
+<p>... and many others.  With a MnemonicAlias definition, the mnemonic is
+remapped simply and directly.  Though MnemonicAlias's can't look at any aspect
+of the instruction (such as the operands) they can depend on global modes (the
+same ones supported by the matcher), through a Requires clause:</p>
+
+<div class="doc_code">
+<pre>
+def : MnemonicAlias&lt;"pushf", "pushfq"&gt;, Requires&lt;[In64BitMode]&gt;;
+def : MnemonicAlias&lt;"pushf", "pushfl"&gt;, Requires&lt;[In32BitMode]&gt;;
+</pre>
+</div>
+
+<p>In this example, the mnemonic gets mapped into different a new one depending
+on the current instruction set.</p>
+
  </div>
  </div>
-<div class="doc_text"><p>To Be Written</p></div>
+
  <!-- _______________________________________________________________________ -->
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
-  <a name="codeemit_bin">Generating Binary Machine Code</a>
+<h4>Instruction Aliases</h4>
+
+<div>
+
+<p>The most general phase of alias processing occurs while matching is
+happening: it provides new forms for the matcher to match along with a specific
+instruction to generate.  An instruction alias has two parts: the string to
+match and the instruction to generate.  For example:
+</p>
+
+<div class="doc_code">
+<pre>
+def : InstAlias&lt;"movsx $src, $dst", (MOVSX16rr8W GR16:$dst, GR8  :$src)&gt;;
+def : InstAlias&lt;"movsx $src, $dst", (MOVSX16rm8W GR16:$dst, i8mem:$src)&gt;;
+def : InstAlias&lt;"movsx $src, $dst", (MOVSX32rr8  GR32:$dst, GR8  :$src)&gt;;
+def : InstAlias&lt;"movsx $src, $dst", (MOVSX32rr16 GR32:$dst, GR16 :$src)&gt;;
+def : InstAlias&lt;"movsx $src, $dst", (MOVSX64rr8  GR64:$dst, GR8  :$src)&gt;;
+def : InstAlias&lt;"movsx $src, $dst", (MOVSX64rr16 GR64:$dst, GR16 :$src)&gt;;
+def : InstAlias&lt;"movsx $src, $dst", (MOVSX64rr32 GR64:$dst, GR32 :$src)&gt;;
+</pre>
+</div>
+
+<p>This shows a powerful example of the instruction aliases, matching the
+same mnemonic in multiple different ways depending on what operands are present
+in the assembly.  The result of instruction aliases can include operands in a
+different order than the destination instruction, and can use an input
+multiple times, for example:</p>
+
+<div class="doc_code">
+<pre>
+def : InstAlias&lt;"clrb $reg", (XOR8rr  GR8 :$reg, GR8 :$reg)&gt;;
+def : InstAlias&lt;"clrw $reg", (XOR16rr GR16:$reg, GR16:$reg)&gt;;
+def : InstAlias&lt;"clrl $reg", (XOR32rr GR32:$reg, GR32:$reg)&gt;;
+def : InstAlias&lt;"clrq $reg", (XOR64rr GR64:$reg, GR64:$reg)&gt;;
+</pre>
+</div>
+
+<p>This example also shows that tied operands are only listed once.  In the X86
+backend, XOR8rr has two input GR8's and one output GR8 (where an input is tied
+to the output).  InstAliases take a flattened operand list without duplicates
+for tied operands.  The result of an instruction alias can also use immediates
+and fixed physical registers which are added as simple immediate operands in the
+result, for example:</p>
+
+<div class="doc_code">
+<pre>
+// Fixed Immediate operand.
+def : InstAlias&lt;"aad", (AAD8i8 10)&gt;;
+
+// Fixed register operand.
+def : InstAlias&lt;"fcomi", (COM_FIr ST1)&gt;;
+
+// Simple alias.
+def : InstAlias&lt;"fcomi $reg", (COM_FIr RST:$reg)&gt;;
+</pre>
+</div>
+
+
+<p>Instruction aliases can also have a Requires clause to make them
+subtarget specific.</p>
+
+<p>If the back-end supports it, the instruction printer can automatically emit
+   the alias rather than what's being aliased. It typically leads to better,
+   more readable code. If it's better to print out what's being aliased, then
+   pass a '0' as the third parameter to the InstAlias definition.</p>
+
  </div>
  
  </div>
  
-<div class="doc_text">
-   <p>For the JIT or <tt>.o</tt> file writer</p>
  </div>
  
  </div>
  
+<!-- ======================================================================= -->
+<h3 id="na_matching">Instruction Matching</h3>
+
+<div><p>To Be Written</p></div>
+
+</div>
  
  <!-- *********************************************************************** -->
  
  <!-- *********************************************************************** -->
-<div class="doc_section">
+<h2>
    <a name="targetimpls">Target-specific Implementation Notes</a>
    <a name="targetimpls">Target-specific Implementation Notes</a>
-</div>
+</h2>
  <!-- *********************************************************************** -->
  
  <!-- *********************************************************************** -->
  
-<div class="doc_text">
+<div>
  
  <p>This section of the document explains features or design decisions that are
  
  <p>This section of the document explains features or design decisions that are
-   specific to the code generator for a particular target.</p>
+   specific to the code generator for a particular target.  First we start
+   with a table that summarizes what features are supported by each target.</p>
+
+<!-- ======================================================================= -->
+<h3>
+  <a name="targetfeatures">Target Feature Matrix</a>
+</h3>
+
+<div>
+
+<p>Note that this table does not include the C backend or Cpp backends, since
+they do not use the target independent code generator infrastructure.  It also
+doesn't list features that are not supported fully by any target yet.  It
+considers a feature to be supported if at least one subtarget supports it.  A
+feature being supported means that it is useful and works for most cases, it
+does not indicate that there are zero known bugs in the implementation.  Here
+is the key:</p>
+
+
+<table border="1" cellspacing="0">
+  <tr>
+    <th>Unknown</th>
+    <th>No support</th>
+    <th>Partial Support</th>
+    <th>Complete Support</th>
+  </tr>
+  <tr>
+    <td class="unknown"></td>
+    <td class="no"></td>
+    <td class="partial"></td>
+    <td class="yes"></td>
+  </tr>
+</table>
+
+<p>Here is the table:</p>
+
+<table width="689" border="1" cellspacing="0">
+<tr><td></td>
+<td colspan="13" align="center" style="background-color:#ffc">Target</td>
+</tr>
+  <tr>
+    <th>Feature</th>
+    <th>ARM</th>
+    <th>Alpha</th>
+    <th>Blackfin</th>
+    <th>CellSPU</th>
+    <th>MBlaze</th>
+    <th>MSP430</th>
+    <th>Mips</th>
+    <th>PTX</th>
+    <th>PowerPC</th>
+    <th>Sparc</th>
+    <th>SystemZ</th>
+    <th>X86</th>
+    <th>XCore</th>
+  </tr>
+
+<tr>
+  <td><a href="#feat_reliable">is generally reliable</a></td>
+  <td class="yes"></td> <!-- ARM -->
+  <td class="unknown"></td> <!-- Alpha -->
+  <td class="no"></td> <!-- Blackfin -->
+  <td class="no"></td> <!-- CellSPU -->
+  <td class="no"></td> <!-- MBlaze -->
+  <td class="unknown"></td> <!-- MSP430 -->
+  <td class="no"></td> <!-- Mips -->
+  <td class="no"></td> <!-- PTX -->
+  <td class="yes"></td> <!-- PowerPC -->
+  <td class="yes"></td> <!-- Sparc -->
+  <td class="unknown"></td> <!-- SystemZ -->
+  <td class="yes"></td> <!-- X86 -->
+  <td class="unknown"></td> <!-- XCore -->
+</tr>
+
+<tr>
+  <td><a href="#feat_asmparser">assembly parser</a></td>
+  <td class="no"></td> <!-- ARM -->
+  <td class="no"></td> <!-- Alpha -->
+  <td class="no"></td> <!-- Blackfin -->
+  <td class="no"></td> <!-- CellSPU -->
+  <td class="yes"></td> <!-- MBlaze -->
+  <td class="no"></td> <!-- MSP430 -->
+  <td class="no"></td> <!-- Mips -->
+  <td class="no"></td> <!-- PTX -->
+  <td class="no"></td> <!-- PowerPC -->
+  <td class="no"></td> <!-- Sparc -->
+  <td class="no"></td> <!-- SystemZ -->
+  <td class="yes"></td> <!-- X86 -->
+  <td class="no"></td> <!-- XCore -->
+</tr>
+
+<tr>
+  <td><a href="#feat_disassembler">disassembler</a></td>
+  <td class="yes"></td> <!-- ARM -->
+  <td class="no"></td> <!-- Alpha -->
+  <td class="no"></td> <!-- Blackfin -->
+  <td class="no"></td> <!-- CellSPU -->
+  <td class="yes"></td> <!-- MBlaze -->
+  <td class="no"></td> <!-- MSP430 -->
+  <td class="no"></td> <!-- Mips -->
+  <td class="no"></td> <!-- PTX -->
+  <td class="no"></td> <!-- PowerPC -->
+  <td class="no"></td> <!-- Sparc -->
+  <td class="no"></td> <!-- SystemZ -->
+  <td class="yes"></td> <!-- X86 -->
+  <td class="no"></td> <!-- XCore -->
+</tr>
+
+<tr>
+  <td><a href="#feat_inlineasm">inline asm</a></td>
+  <td class="yes"></td> <!-- ARM -->
+  <td class="unknown"></td> <!-- Alpha -->
+  <td class="yes"></td> <!-- Blackfin -->
+  <td class="no"></td> <!-- CellSPU -->
+  <td class="yes"></td> <!-- MBlaze -->
+  <td class="unknown"></td> <!-- MSP430 -->
+  <td class="no"></td> <!-- Mips -->
+  <td class="unknown"></td> <!-- PTX -->
+  <td class="yes"></td> <!-- PowerPC -->
+  <td class="unknown"></td> <!-- Sparc -->
+  <td class="unknown"></td> <!-- SystemZ -->
+  <td class="yes"></td> <!-- X86 -->
+  <td class="unknown"></td> <!-- XCore -->
+</tr>
+
+<tr>
+  <td><a href="#feat_jit">jit</a></td>
+  <td class="partial"><a href="#feat_jit_arm">*</a></td> <!-- ARM -->
+  <td class="no"></td> <!-- Alpha -->
+  <td class="no"></td> <!-- Blackfin -->
+  <td class="no"></td> <!-- CellSPU -->
+  <td class="no"></td> <!-- MBlaze -->
+  <td class="unknown"></td> <!-- MSP430 -->
+  <td class="no"></td> <!-- Mips -->
+  <td class="unknown"></td> <!-- PTX -->
+  <td class="yes"></td> <!-- PowerPC -->
+  <td class="unknown"></td> <!-- Sparc -->
+  <td class="unknown"></td> <!-- SystemZ -->
+  <td class="yes"></td> <!-- X86 -->
+  <td class="unknown"></td> <!-- XCore -->
+</tr>
+
+<tr>
+  <td><a href="#feat_objectwrite">.o&nbsp;file writing</a></td>
+  <td class="no"></td> <!-- ARM -->
+  <td class="no"></td> <!-- Alpha -->
+  <td class="no"></td> <!-- Blackfin -->
+  <td class="no"></td> <!-- CellSPU -->
+  <td class="yes"></td> <!-- MBlaze -->
+  <td class="no"></td> <!-- MSP430 -->
+  <td class="no"></td> <!-- Mips -->
+  <td class="no"></td> <!-- PTX -->
+  <td class="no"></td> <!-- PowerPC -->
+  <td class="no"></td> <!-- Sparc -->
+  <td class="no"></td> <!-- SystemZ -->
+  <td class="yes"></td> <!-- X86 -->
+  <td class="no"></td> <!-- XCore -->
+</tr>
+
+<tr>
+  <td><a href="#feat_tailcall">tail calls</a></td>
+  <td class="yes"></td> <!-- ARM -->
+  <td class="unknown"></td> <!-- Alpha -->
+  <td class="no"></td> <!-- Blackfin -->
+  <td class="no"></td> <!-- CellSPU -->
+  <td class="no"></td> <!-- MBlaze -->
+  <td class="unknown"></td> <!-- MSP430 -->
+  <td class="no"></td> <!-- Mips -->
+  <td class="unknown"></td> <!-- PTX -->
+  <td class="yes"></td> <!-- PowerPC -->
+  <td class="unknown"></td> <!-- Sparc -->
+  <td class="unknown"></td> <!-- SystemZ -->
+  <td class="yes"></td> <!-- X86 -->
+  <td class="unknown"></td> <!-- XCore -->
+</tr>
+
+
+</table>
+
+<!-- _______________________________________________________________________ -->
+<h4 id="feat_reliable">Is Generally Reliable</h4>
+
+<div>
+<p>This box indicates whether the target is considered to be production quality.
+This indicates that the target has been used as a static compiler to
+compile large amounts of code by a variety of different people and is in
+continuous use.</p>
+</div>
+
+<!-- _______________________________________________________________________ -->
+<h4 id="feat_asmparser">Assembly Parser</h4>
+
+<div>
+<p>This box indicates whether the target supports parsing target specific .s
+files by implementing the MCAsmParser interface.  This is required for llvm-mc
+to be able to act as a native assembler and is required for inline assembly
+support in the native .o file writer.</p>
+
+</div>
+
+
+<!-- _______________________________________________________________________ -->
+<h4 id="feat_disassembler">Disassembler</h4>
+
+<div>
+<p>This box indicates whether the target supports the MCDisassembler API for
+disassembling machine opcode bytes into MCInst's.</p>
+
+</div>
+
+<!-- _______________________________________________________________________ -->
+<h4 id="feat_inlineasm">Inline Asm</h4>
+
+<div>
+<p>This box indicates whether the target supports most popular inline assembly
+constraints and modifiers.</p>
+
+</div>
+
+<!-- _______________________________________________________________________ -->
+<h4 id="feat_jit">JIT Support</h4>
+
+<div>
+<p>This box indicates whether the target supports the JIT compiler through
+the ExecutionEngine interface.</p>
+
+<p id="feat_jit_arm">The ARM backend has basic support for integer code
+in ARM codegen mode, but lacks NEON and full Thumb support.</p>
+
+</div>
+
+<!-- _______________________________________________________________________ -->
+<h4 id="feat_objectwrite">.o File Writing</h4>
+
+<div>
+
+<p>This box indicates whether the target supports writing .o files (e.g. MachO,
+ELF, and/or COFF) files directly from the target.  Note that the target also
+must include an assembly parser and general inline assembly support for full
+inline assembly support in the .o writer.</p>
+
+<p>Targets that don't support this feature can obviously still write out .o
+files, they just rely on having an external assembler to translate from a .s
+file to a .o file (as is the case for many C compilers).</p>
+
+</div>
+
+<!-- _______________________________________________________________________ -->
+<h4 id="feat_tailcall">Tail Calls</h4>
+
+<div>
+
+<p>This box indicates whether the target supports guaranteed tail calls.  These
+are calls marked "<a href="LangRef.html#i_call">tail</a>" and use the fastcc
+calling convention.  Please see the <a href="#tailcallopt">tail call section
+more more details</a>.</p>
+
+</div>
  
  </div>
  
  <!-- ======================================================================= -->
  
  </div>
  
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="tailcallopt">Tail call optimization</a>
    <a name="tailcallopt">Tail call optimization</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>Tail call optimization, callee reusing the stack of the caller, is currently
     supported on x86/x86-64 and PowerPC. It is performed if:</p>
  
  <p>Tail call optimization, callee reusing the stack of the caller, is currently
     supported on x86/x86-64 and PowerPC. It is performed if:</p>
@@ -1735,11 +2493,11 @@ define fastcc i32 @tailcaller(i32 %in1, i32 %in2) {
  
  </div>
  <!-- ======================================================================= -->
  
  </div>
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="sibcallopt">Sibling call optimization</a>
    <a name="sibcallopt">Sibling call optimization</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>Sibling call optimization is a restricted form of tail call optimization.
     Unlike tail call optimization described in the previous section, it can be
  
  <p>Sibling call optimization is a restricted form of tail call optimization.
     Unlike tail call optimization described in the previous section, it can be
@@ -1779,24 +2537,22 @@ entry:
  
  </div>
  <!-- ======================================================================= -->
  
  </div>
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="x86">The X86 backend</a>
    <a name="x86">The X86 backend</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The X86 code generator lives in the <tt>lib/Target/X86</tt> directory.  This
     code generator is capable of targeting a variety of x86-32 and x86-64
     processors, and includes support for ISA extensions such as MMX and SSE.</p>
  
  
  <p>The X86 code generator lives in the <tt>lib/Target/X86</tt> directory.  This
     code generator is capable of targeting a variety of x86-32 and x86-64
     processors, and includes support for ISA extensions such as MMX and SSE.</p>
  
-</div>
-
  <!-- _______________________________________________________________________ -->
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="x86_tt">X86 Target Triples supported</a>
    <a name="x86_tt">X86 Target Triples supported</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The following are the known target triples that are supported by the X86
     backend.  This is not an exhaustive list, and it would be useful to add those
  
  <p>The following are the known target triples that are supported by the X86
     backend.  This is not an exhaustive list, and it would be useful to add those
@@ -1821,31 +2577,34 @@ entry:
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="x86_cc">X86 Calling Conventions supported</a>
    <a name="x86_cc">X86 Calling Conventions supported</a>
-</div>
+</h4>
  
  
  
  
-<div class="doc_text">
+<div>
  
  <p>The following target-specific calling conventions are known to backend:</p>
  
  <ul>
  
  <p>The following target-specific calling conventions are known to backend:</p>
  
  <ul>
-  <li><b>x86_StdCall</b> &mdash; stdcall calling convention seen on Microsoft
-      Windows platform (CC ID = 64).</li>
-
-  <li><b>x86_FastCall</b> &mdash; fastcall calling convention seen on Microsoft
-      Windows platform (CC ID = 65).</li>
+<li><b>x86_StdCall</b> &mdash; stdcall calling convention seen on Microsoft
+    Windows platform (CC ID = 64).</li>
+<li><b>x86_FastCall</b> &mdash; fastcall calling convention seen on Microsoft
+    Windows platform (CC ID = 65).</li>
+<li><b>x86_ThisCall</b> &mdash; Similar to X86_StdCall. Passes first argument
+    in ECX,  others via stack. Callee is responsible for stack cleaning. This
+    convention is used by MSVC by default for methods in its ABI
+    (CC ID = 70).</li>
  </ul>
  
  </div>
  
  <!-- _______________________________________________________________________ -->
  </ul>
  
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="x86_memory">Representing X86 addressing modes in MachineInstrs</a>
    <a name="x86_memory">Representing X86 addressing modes in MachineInstrs</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The x86 has a very flexible way of accessing memory.  It is capable of
     forming memory addresses of the following expression directly in integer
  
  <p>The x86 has a very flexible way of accessing memory.  It is capable of
     forming memory addresses of the following expression directly in integer
@@ -1878,13 +2637,13 @@ OperandTy: VirtReg, | VirtReg, UnsImm, VirtReg,   SignExtImm  PhysReg
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="x86_memory">X86 address spaces supported</a>
    <a name="x86_memory">X86 address spaces supported</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  
-<p>x86 has an experimental feature which provides
+<p>x86 has a feature which provides
     the ability to perform loads and stores to different address spaces
     via the x86 segment registers.  A segment override prefix byte on an
     instruction causes the instruction's memory access to go to the specified
     the ability to perform loads and stores to different address spaces
     via the x86 segment registers.  A segment override prefix byte on an
     instruction causes the instruction's memory access to go to the specified
@@ -1923,11 +2682,11 @@ OperandTy: VirtReg, | VirtReg, UnsImm, VirtReg,   SignExtImm  PhysReg
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="x86_names">Instruction naming</a>
    <a name="x86_names">Instruction naming</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>An instruction name consists of the base name, a default operand size, and a
     a character per operand with an optional special size. For example:</p>
  
  <p>An instruction name consists of the base name, a default operand size, and a
     a character per operand with an optional special size. For example:</p>
@@ -1943,25 +2702,25 @@ MOVSX32rm16 -&gt; movsx, 32-bit register, 16-bit memory
  
  </div>
  
  
  </div>
  
+</div>
+
  <!-- ======================================================================= -->
  <!-- ======================================================================= -->
-<div class="doc_subsection">
+<h3>
    <a name="ppc">The PowerPC backend</a>
    <a name="ppc">The PowerPC backend</a>
-</div>
+</h3>
  
  
-<div class="doc_text">
+<div>
  
  <p>The PowerPC code generator lives in the lib/Target/PowerPC directory.  The
     code generation is retargetable to several variations or <i>subtargets</i> of
     the PowerPC ISA; including ppc32, ppc64 and altivec.</p>
  
  
  <p>The PowerPC code generator lives in the lib/Target/PowerPC directory.  The
     code generation is retargetable to several variations or <i>subtargets</i> of
     the PowerPC ISA; including ppc32, ppc64 and altivec.</p>
  
-</div>
-
  <!-- _______________________________________________________________________ -->
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="ppc_abi">LLVM PowerPC ABI</a>
    <a name="ppc_abi">LLVM PowerPC ABI</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>LLVM follows the AIX PowerPC ABI, with two deviations. LLVM uses a PC
     relative (PIC) or static addressing for accessing global values, so no TOC
  
  <p>LLVM follows the AIX PowerPC ABI, with two deviations. LLVM uses a PC
     relative (PIC) or static addressing for accessing global values, so no TOC
@@ -1977,11 +2736,11 @@ MOVSX32rm16 -&gt; movsx, 32-bit register, 16-bit memory
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="ppc_frame">Frame Layout</a>
    <a name="ppc_frame">Frame Layout</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The size of a PowerPC frame is usually fixed for the duration of a
     function's invocation.  Since the frame is fixed size, all references
  
  <p>The size of a PowerPC frame is usually fixed for the duration of a
     function's invocation.  Since the frame is fixed size, all references
@@ -2124,11 +2883,11 @@ MOVSX32rm16 -&gt; movsx, 32-bit register, 16-bit memory
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="ppc_prolog">Prolog/Epilog</a>
    <a name="ppc_prolog">Prolog/Epilog</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p>The llvm prolog and epilog are the same as described in the PowerPC ABI, with
     the following exceptions.  Callee saved registers are spilled after the frame
  
  <p>The llvm prolog and epilog are the same as described in the PowerPC ABI, with
     the following exceptions.  Callee saved registers are spilled after the frame
@@ -2141,16 +2900,83 @@ MOVSX32rm16 -&gt; movsx, 32-bit register, 16-bit memory
  </div>
  
  <!-- _______________________________________________________________________ -->
  </div>
  
  <!-- _______________________________________________________________________ -->
-<div class="doc_subsubsection">
+<h4>
    <a name="ppc_dynamic">Dynamic Allocation</a>
    <a name="ppc_dynamic">Dynamic Allocation</a>
-</div>
+</h4>
  
  
-<div class="doc_text">
+<div>
  
  <p><i>TODO - More to come.</i></p>
  
  </div>
  
  
  <p><i>TODO - More to come.</i></p>
  
  </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>
  
  <!-- *********************************************************************** -->
  <hr>
  
  <!-- *********************************************************************** -->
  <hr>
@@ -2161,7 +2987,7 @@ MOVSX32rm16 -&gt; movsx, 32-bit register, 16-bit memory
    src="http://www.w3.org/Icons/valid-html401-blue" alt="Valid HTML 4.01"></a>
  
    <a href="mailto:sabre@nondot.org">Chris Lattner</a><br>
    src="http://www.w3.org/Icons/valid-html401-blue" alt="Valid HTML 4.01"></a>
  
    <a href="mailto:sabre@nondot.org">Chris Lattner</a><br>
-  <a href="http://llvm.org">The LLVM Compiler Infrastructure</a><br>
+  <a href="http://llvm.org/">The LLVM Compiler Infrastructure</a><br>
    Last modified: $Date$
  </address>
  
    Last modified: $Date$
  </address>