RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-12-03 17:02:03 +00:00
Code Issues Actions Packages Projects Releases Wiki Activity

Remove PTX->NVPTX in CodeGenerator document and update its text.

Browse Source 
llvm-svn: 172235
This commit is contained in:
Justin Holewinski 2013-01-11 18:47:10 +00:00
parent 42297e1413
commit 4744632031
1 changed files with 36 additions and 47 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

83
docs/CodeGenerator.rst
View File

@ -1775,7 +1775,7 @@ Here is the table:
:raw-html:`<th>MBlaze</th>`
:raw-html:`<th>MSP430</th>`
:raw-html:`<th>Mips</th>`
:raw-html:`<th>PTX</th>`
:raw-html:`<th>NVPTX</th>`
:raw-html:`<th>PowerPC</th>`
:raw-html:`<th>Sparc</th>`
:raw-html:`<th>X86</th>`
@ -1789,7 +1789,7 @@ Here is the table:
:raw-html:`<td class="no"></td> <!-- MBlaze -->`
:raw-html:`<td class="unknown"></td> <!-- MSP430 -->`
:raw-html:`<td class="yes"></td> <!-- Mips -->`
:raw-html:`<td class="yes"></td> <!-- PTX -->`
:raw-html:`<td class="yes"></td> <!-- NVPTX -->`
:raw-html:`<td class="yes"></td> <!-- PowerPC -->`
:raw-html:`<td class="yes"></td> <!-- Sparc -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
@ -1803,7 +1803,7 @@ Here is the table:
:raw-html:`<td class="yes"></td> <!-- MBlaze -->`
:raw-html:`<td class="no"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
:raw-html:`<td class="no"></td> <!-- PTX -->`
:raw-html:`<td class="no"></td> <!-- NVPTX -->`
:raw-html:`<td class="no"></td> <!-- PowerPC -->`
:raw-html:`<td class="no"></td> <!-- Sparc -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
@ -1817,7 +1817,7 @@ Here is the table:
:raw-html:`<td class="yes"></td> <!-- MBlaze -->`
:raw-html:`<td class="no"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
:raw-html:`<td class="na"></td> <!-- PTX -->`
:raw-html:`<td class="na"></td> <!-- NVPTX -->`
:raw-html:`<td class="no"></td> <!-- PowerPC -->`
:raw-html:`<td class="no"></td> <!-- Sparc -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
@ -1831,7 +1831,7 @@ Here is the table:
:raw-html:`<td class="yes"></td> <!-- MBlaze -->`
:raw-html:`<td class="unknown"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
:raw-html:`<td class="yes"></td> <!-- PTX -->`
:raw-html:`<td class="yes"></td> <!-- NVPTX -->`
:raw-html:`<td class="yes"></td> <!-- PowerPC -->`
:raw-html:`<td class="unknown"></td> <!-- Sparc -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
@ -1845,7 +1845,7 @@ Here is the table:
:raw-html:`<td class="no"></td> <!-- MBlaze -->`
:raw-html:`<td class="unknown"></td> <!-- MSP430 -->`
:raw-html:`<td class="yes"></td> <!-- Mips -->`
:raw-html:`<td class="na"></td> <!-- PTX -->`
:raw-html:`<td class="na"></td> <!-- NVPTX -->`
:raw-html:`<td class="yes"></td> <!-- PowerPC -->`
:raw-html:`<td class="unknown"></td> <!-- Sparc -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
@ -1859,7 +1859,7 @@ Here is the table:
:raw-html:`<td class="yes"></td> <!-- MBlaze -->`
:raw-html:`<td class="no"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
:raw-html:`<td class="na"></td> <!-- PTX -->`
:raw-html:`<td class="na"></td> <!-- NVPTX -->`
:raw-html:`<td class="no"></td> <!-- PowerPC -->`
:raw-html:`<td class="no"></td> <!-- Sparc -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
@ -1873,7 +1873,7 @@ Here is the table:
:raw-html:`<td class="no"></td> <!-- MBlaze -->`
:raw-html:`<td class="unknown"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
:raw-html:`<td class="no"></td> <!-- PTX -->`
:raw-html:`<td class="no"></td> <!-- NVPTX -->`
:raw-html:`<td class="yes"></td> <!-- PowerPC -->`
:raw-html:`<td class="unknown"></td> <!-- Sparc -->`
:raw-html:`<td class="yes"></td> <!-- X86 -->`
@ -1887,7 +1887,7 @@ Here is the table:
:raw-html:`<td class="no"></td> <!-- MBlaze -->`
:raw-html:`<td class="no"></td> <!-- MSP430 -->`
:raw-html:`<td class="no"></td> <!-- Mips -->`
:raw-html:`<td class="no"></td> <!-- PTX -->`
:raw-html:`<td class="no"></td> <!-- NVPTX -->`
:raw-html:`<td class="no"></td> <!-- PowerPC -->`
:raw-html:`<td class="no"></td> <!-- Sparc -->`
:raw-html:`<td class="partial"><a href="#feat_segstacks_x86">*</a></td> <!-- X86 -->`
@ -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:`<th>Description</th>`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td>``double``</td>`
:raw-html:`<td align="left">If enabled, the map_f64_to_f32 directive is disabled in the PTX output, allowing native double-precision arithmetic</td>`
:raw-html:`<td>sm_20</td>`
:raw-html:`<td align="left">Set shader model/compute capability to 2.0</td>`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td>``no-fma``</td>`
:raw-html:`<td align="left">Disable generation of Fused-Multiply Add instructions, which may be beneficial for some devices</td>`
:raw-html:`<td>sm_21</td>`
:raw-html:`<td align="left">Set shader model/compute capability to 2.1</td>`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td>``smxy / computexy``</td>`
:raw-html:`<td align="left">Set shader model/compute capability to x.y, e.g. sm20 or compute13</td>`
:raw-html:`<td>sm_30</td>`
:raw-html:`<td align="left">Set shader model/compute capability to 3.0</td>`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td>sm_35</td>`
:raw-html:`<td align="left">Set shader model/compute capability to 3.5</td>`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td>ptx30</td>`
:raw-html:`<td align="left">Target PTX 3.0</td>`
:raw-html:`</tr>`
:raw-html:`<tr>`
:raw-html:`<td>ptx31</td>`
:raw-html:`<td align="left">Target PTX 3.1</td>`
:raw-html:`</tr>`
:raw-html:`</table>`
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