Durgadoss R 1b01064faa
[NVPTX] Add TMA bulk tensor copy intrinsics (#96083)
This patch adds NVVM intrinsics and NVPTX codegen for:
* cp.async.bulk.tensor.S2G.1D -> 5D variants, supporting both Tile and
   Im2Col modes. These intrinsics optionally support cache_hints as
   indicated by the boolean flag argument.
* cp.async.bulk.tensor.G2S.1D -> 5D variants, with support for both Tile
  and Im2Col modes. The Im2Col variants have an extra set of offsets as
  parameters. These intrinsics optionally support multicast and cache_hints,
  as indicated by the boolean arguments at the end of the intrinsics.
* The backend looks through these flag arguments and lowers to the
   appropriate PTX instruction.
* Lit tests are added for all combinations of these intrinsics in
  cp-async-bulk-tensor-g2s/s2g.ll.
* The generated PTX is verified with a 12.3 ptxas executable.
* Added docs for these intrinsics in NVPTXUsage.rst file.
* PTX Spec reference:
  https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-async-bulk-tensor

Signed-off-by: Durgadoss R <durgadossr@nvidia.com>
2024-11-07 15:21:53 +05:30
..
2024-05-19 22:09:46 +05:30
2024-07-17 21:57:52 -07:00
2024-01-09 11:06:33 +08:00

LLVM Documentation
==================

LLVM's documentation is written in reStructuredText, a lightweight
plaintext markup language (file extension `.rst`). While the
reStructuredText documentation should be quite readable in source form, it
is mostly meant to be processed by the Sphinx documentation generation
system to create HTML pages which are hosted on <https://llvm.org/docs/> and
updated after every commit. Manpage output is also supported, see below.

If you instead would like to generate and view the HTML locally, install
Sphinx <http://sphinx-doc.org/> and then do:

    cd <build-dir>
    cmake -DLLVM_ENABLE_SPHINX=true -DSPHINX_OUTPUT_HTML=true <src-dir>
    make -j3 docs-llvm-html
    $BROWSER <build-dir>/docs/html/index.html

The mapping between reStructuredText files and generated documentation is
`docs/Foo.rst` <-> `<build-dir>/docs//html/Foo.html` <-> `https://llvm.org/docs/Foo.html`.

If you are interested in writing new documentation, you will want to read
`SphinxQuickstartTemplate.rst` which will get you writing documentation
very fast and includes examples of the most important reStructuredText
markup syntax.

Manpage Output
===============

Building the manpages is similar to building the HTML documentation. The
primary difference is to use the `man` makefile target, instead of the
default (which is `html`). Sphinx then produces the man pages in the
directory `<build-dir>/docs/man/`.

    cd <build-dir>
    cmake -DLLVM_ENABLE_SPHINX=true -DSPHINX_OUTPUT_MAN=true <src-dir>
    make -j3 docs-llvm-man
    man -l <build-dir>/docs/man/FileCheck.1

The correspondence between .rst files and man pages is
`docs/CommandGuide/Foo.rst` <-> `<build-dir>/docs//man/Foo.1`.
These .rst files are also included during HTML generation so they are also
viewable online (as noted above) at e.g.
`https://llvm.org/docs/CommandGuide/Foo.html`.

Checking links
==============

The reachability of external links in the documentation can be checked by
running:

    cd llvm/docs/
    sphinx-build -b linkcheck . _build/lintcheck/
    # report will be generated in _build/lintcheck/output.txt

Doxygen page Output
==============

Install doxygen <https://www.doxygen.nl/download.html> and dot2tex <https://dot2tex.readthedocs.io/en/latest>.

    cd <build-dir>
    cmake -DLLVM_ENABLE_DOXYGEN=On <llvm-top-src-dir>
    make doxygen-llvm # for LLVM docs
    make doxygen-clang # for clang docs

It will generate html in

    <build-dir>/docs/doxygen/html # for LLVM docs
    <build-dir>/tools/clang/docs/doxygen/html # for clang docs