diff --git a/docs/CodeGenerator.rst b/docs/CodeGenerator.rst
index 515c57631e3..0393d317bb8 100644
--- a/docs/CodeGenerator.rst
+++ b/docs/CodeGenerator.rst
@@ -1775,7 +1775,7 @@ Here is the table:
:raw-html:`
MBlaze | `
:raw-html:`MSP430 | `
:raw-html:`Mips | `
-:raw-html:`PTX | `
+:raw-html:`NVPTX | `
:raw-html:`PowerPC | `
:raw-html:`Sparc | `
:raw-html:`X86 | `
@@ -1789,7 +1789,7 @@ Here is the table:
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
-:raw-html:` | `
+:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
@@ -1803,7 +1803,7 @@ Here is the table:
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
-:raw-html:` | `
+:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
@@ -1817,7 +1817,7 @@ Here is the table:
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
-:raw-html:` | `
+:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
@@ -1831,7 +1831,7 @@ Here is the table:
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
-:raw-html:` | `
+:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
@@ -1845,7 +1845,7 @@ Here is the table:
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
-:raw-html:` | `
+:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
@@ -1859,7 +1859,7 @@ Here is the table:
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
-:raw-html:` | `
+:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
@@ -1873,7 +1873,7 @@ Here is the table:
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
-:raw-html:` | `
+:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
@@ -1887,7 +1887,7 @@ Here is the table:
:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
-:raw-html:` | `
+:raw-html:` | `
:raw-html:` | `
:raw-html:` | `
:raw-html:`* | `
@@ -2369,17 +2369,17 @@ Dynamic Allocation
TODO - More to come.
-The PTX backend
----------------
+The NVPTX backend
+-----------------
-The PTX code generator lives in the lib/Target/PTX directory. It is currently a
-work-in-progress, but already supports most of the code generation functionality
-needed to generate correct PTX kernels for CUDA devices.
+The NVPTX code generator under lib/Target/NVPTX is an open-source version of
+the NVIDIA NVPTX code generator for LLVM. It is contributed by NVIDIA and is
+a port of the code generator used in the CUDA compiler (nvcc). It targets the
+PTX 3.0/3.1 ISA and can target any compute capability greater than or equal to
+2.0 (Fermi).
-The code generator can target PTX 2.0+, and shader model 1.0+. The PTX ISA
-Reference Manual is used as the primary source of ISA information, though an
-effort is made to make the output of the code generator match the output of the
-NVidia nvcc compiler, whenever possible.
+This target is of production quality and should be completely compatible with
+the official NVIDIA toolchain.
Code Generator Options:
@@ -2389,39 +2389,28 @@ Code Generator Options:
:raw-html:`Description | `
:raw-html:``
:raw-html:``
-:raw-html:```double`` | `
-:raw-html:`If enabled, the map_f64_to_f32 directive is disabled in the PTX output, allowing native double-precision arithmetic | `
+:raw-html:`sm_20 | `
+:raw-html:`Set shader model/compute capability to 2.0 | `
:raw-html:`
`
:raw-html:``
-:raw-html:```no-fma`` | `
-:raw-html:`Disable generation of Fused-Multiply Add instructions, which may be beneficial for some devices | `
+:raw-html:`sm_21 | `
+:raw-html:`Set shader model/compute capability to 2.1 | `
:raw-html:`
`
:raw-html:``
-:raw-html:```smxy / computexy`` | `
-:raw-html:`Set shader model/compute capability to x.y, e.g. sm20 or compute13 | `
+:raw-html:`sm_30 | `
+:raw-html:`Set shader model/compute capability to 3.0 | `
+:raw-html:`
`
+:raw-html:``
+:raw-html:`sm_35 | `
+:raw-html:`Set shader model/compute capability to 3.5 | `
+:raw-html:`
`
+:raw-html:``
+:raw-html:`ptx30 | `
+:raw-html:`Target PTX 3.0 | `
+:raw-html:`
`
+:raw-html:``
+:raw-html:`ptx31 | `
+:raw-html:`Target PTX 3.1 | `
:raw-html:`
`
:raw-html:``
-Working:
-
-* Arithmetic instruction selection (including combo FMA)
-
-* Bitwise instruction selection
-
-* Control-flow instruction selection
-
-* Function calls (only on SM 2.0+ and no return arguments)
-
-* Addresses spaces (0 = global, 1 = constant, 2 = local, 4 = shared)
-
-* Thread synchronization (bar.sync)
-
-* Special register reads ([N]TID, [N]CTAID, PMx, CLOCK, etc.)
-
-In Progress:
-
-* Robust call instruction selection
-
-* Stack frame allocation
-
-* Device-specific instruction scheduling optimizations