From 9c0d0f55da1724d303d4b630101b6444b20a4223 Mon Sep 17 00:00:00 2001 From: Justin Holewinski Date: Fri, 11 Jan 2013 18:47:10 +0000 Subject: Remove PTX->NVPTX in CodeGenerator document and update its text. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@172235 91177308-0d34-0410-b5e6-96231b3b80d8 --- docs/CodeGenerator.rst | 83 ++++++++++++++++++++++---------------------------- 1 file changed, 36 insertions(+), 47 deletions(-) (limited to 'docs/CodeGenerator.rst') diff --git a/docs/CodeGenerator.rst b/docs/CodeGenerator.rst index 515c57631e..0393d317bb 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:`sm_21` +:raw-html:`Set shader model/compute capability to 2.1` +:raw-html:`` +:raw-html:`` +:raw-html:`sm_30` +:raw-html:`Set shader model/compute capability to 3.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_35` +:raw-html:`Set shader model/compute capability to 3.5` :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:`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 -- cgit v1.2.3