[Doc] better documention

docs/atomic.rst

@@ -0,0 +1,89 @@
+.. _atomic:
+
+Atomic operations
+=================
+
+In taichi, ``x[i] += 1`` is atomic but ``x[i] = x[i] + 1`` is not.
+
+For example, to perform a reduction:
+::
+    @ti.kernel
+    def sum():
+        for i in x:
+            result[None] += x[i]
+
+Or use the function ``ti.atomic_add``, which is equivalent:
+::
+    @ti.kernel
+    def sum():
+        for i in x:
+            ti.atomic_add(result[None], x[i])
+
+See https://en.wikipedia.org/wiki/Fetch-and-add for more details.
+
+
+.. note::
+    Support of atomic operation on each backends:
+
+    +------+-----------+-----------+---------+
+    | type | CPU/CUDA  | OpenGL    | Metal   |
+    +======+===========+===========+=========+
+    | i32  |    OK     |    OK     |   OK    |
+    +------+-----------+-----------+---------+
+    | f32  |    OK     |    OK     |   OK    |
+    +------+-----------+-----------+---------+
+    | i64  |    OK     |   EXT     |  MISS   |
+    +------+-----------+-----------+---------+
+    | f64  |    OK     |   EXT     |  MISS   |
+    +------+-----------+-----------+---------+
+
+    (OK=supported, EXT=require extension, MISS=not supported)
+
+
+.. function:: ti.atomic_add(x, y)
+
+    This is equivalent to ``x += y``.
+
+    :parameter x: (Expr, lvalue) the LHS of addition, also where the result is saved
+    :parameter y: (Expr) the RHS of addition
+    :return: (Expr) the original value stored in ``x``
+
+    e.g.:
+    ::
+        x = 3
+        y = 4
+        z = ti.atomic_add(x, y)
+        # now x = 7, y = 4, z = 3
+
+
+.. function:: ti.atomic_sub(x, y)
+
+    This is equivalent to ``x -= y``.
+
+
+.. function:: ti.atomic_max(x, y)
+
+    e.g.:
+    ::
+        x = 3
+        y = 4
+        z = ti.atomic_max(x, y)
+        # now x = 4, y = 4, z = 3
+
+
+.. function:: ti.atomic_min(x, y)
+
+
+.. function:: ti.atomic_or(x, y)
+
+    This is equivalent to ``x |= y``.
+
+
+.. function:: ti.atomic_and(x, y)
+
+    This is equivalent to ``x &= y``.
+
+
+.. function:: ti.atomic_xor(x, y)
+
+    This is equivalent to ``x ^= y``.

docs/compilation.rst

@@ -101,7 +101,7 @@ which allows a series of further IR passes to happen, such as
  - Atomic operation demotion
 
 The just-in-time (JIT) compilation engine
----------------------------------------
+-----------------------------------------
 
 Finally, the optimized SSA IR is fed into the LLVM IR codegen, and LLVM JIT generates high-performance executable CPU/GPU programs.
 

docs/cpp_style.rst

@@ -28,3 +28,7 @@ Don'ts
 - ``NULL``, use ``nullptr`` instead.
 - ``using namespace std;`` in global scope.
 - ``typedef``. Use ``using`` instead.
+
+Automatic code formatting
+--------------------------------------------------------------------------------
+- Please run ``ti format``

docs/dev_install.rst

@@ -11,6 +11,9 @@ For precise build instructions on Windows, please check out `appveyor.yml <https
 
 Note that on Linux/OS X, ``clang`` is the only supported compiler for compiling the Taichi compiler. On Windows only MSVC supported.
 
+Installing Depedencies
+---------------------------------------------
+
 - Make sure you are using Python 3.6/3.7/3.8
 - Execute
 
@@ -19,39 +22,62 @@ Note that on Linux/OS X, ``clang`` is the only supported compiler for compiling
     python3 -m pip install --user setuptools astpretty astor pytest opencv-python pybind11
     python3 -m pip install --user Pillow numpy scipy GitPython yapf colorama psutil autograd
 
-- (If on Ubuntu) Execute ``sudo apt install libtinfo-dev clang-8``. ``clang-7`` should work as well.
-- Make sure you have LLVM 8.0.1 built from scratch (`Download <https://github.com/llvm/llvm-project/releases/download/llvmorg-8.0.1/llvm-8.0.1.src.tar.xz>`_). To do so, download and unzip the llvm source, move to the llvm folder, and execute
+* (If on Ubuntu) Execute ``sudo apt install libtinfo-dev clang-8`` (or ``clang-7`` should work as well).
+
+* (If on Arch Linux) Execute
 
   .. code-block:: bash
 
+    wget https://archive.archlinux.org/packages/c/clang/clang-8.0.1-1-x86_64.pkg.tar.xz
+    sudo pacman -Qp clang-8.0.1-1-x86_64.pkg.tar.xz
+
+  .. warning::
+    If you have installed ``clang`` (9.0.1) before, this command will overrides the existing ``clang``.
+    If you don't want to break up depedencies, please build from scratch and install it in ``/opt``. Then add ``/opt/clang/bin`` to your ``$PATH``.
+
+
+- Make sure you have LLVM 8.0.1 built from scratch. To do so:
+
+  .. code-block:: bash
+
+    wget https://github.com/llvm/llvm-project/releases/download/llvmorg-8.0.1/llvm-8.0.1.src.tar.xz
+    tar xvJf llvm-8.0.1.src.tar.xz
+    cd llvm-8.0.1.src
     mkdir build
     cd build
     cmake .. -DLLVM_ENABLE_RTTI:BOOL=ON -DBUILD_SHARED_LIBS:BOOL=OFF -DCMAKE_BUILD_TYPE=Release -DLLVM_TARGETS_TO_BUILD="X86;NVPTX" -DLLVM_ENABLE_ASSERTIONS=ON
     # If you are building on NVIDIA Jetson TX2, use -DLLVM_TARGETS_TO_BUILD="ARM;NVPTX"
     make -j 8
     sudo make install
 
-- Clone the taichi repo, and then
+Setting up Taichi for develop
+---------------------------------------------
+
+- Clone the taichi repo, and build:
 
   .. code-block:: bash
 
+    git clone https://github.com/taichi-dev/taichi --depth=1 --branch=master
+    gsm update --init --recursive --depth=1
     cd taichi
     mkdir build
     cd build
     cmake ..
-    # if you are building with CUDA, say, 10.0, then please use "cmake .. -DCUDA_VERSION=10.0 -DTI_WITH_CUDA:BOOL=True"
+    # if you are building with CUDA 10.0, use this:
+    # cmake .. -DCUDA_VERSION=10.0 -DTI_WITH_CUDA:BOOL=True
     make -j 8
 
-- Add the following to your ``~/.bashrc`` (or ``~/.zshrc`` if you use ``zsh``)
+- Add the following codes to your ``~/.bashrc``:
 
   .. code-block:: bash
 
     export TAICHI_REPO_DIR=/home/XXX/taichi  # Path to your taichi repository
     export PYTHONPATH=$TAICHI_REPO_DIR/python/:$PYTHONPATH
     export PATH=$TAICHI_REPO_DIR/bin/:$PATH
+    # export PATH=/opt/llvm/bin:$PATH # Uncomment if your llvm-8 or clang-8 is in /opt
 
 - Execute ``source ~/.bashrc`` to reload shell config
-- Execute ``ti test`` to run all the tests. It may take up to 5 minutes to run all tests. (On Windows the ``ti`` command should be replaced by ``python -m taichi``)
+- Execute ``ti test`` to run all the tests. It may take up to 5 minutes to run all tests. (On Windows, please execute ``python3 -m taichi test`` instead)
 - Check out ``examples`` for runnable examples. Run them with ``python3``.
 
 

docs/global_settings.rst

@@ -3,4 +3,8 @@ Global Settings
 
 - Restart the Taichi runtime system (clear memory, destroy all variables and kernels): ``ti.reset()``
 - Eliminate verbose outputs: ``ti.get_runtime().set_verbose(False)``
-- To specify which GPU to use: ``export CUDA_VISIBLE_DEVICES=0``
+- Not to trigger GDB when crashes: ``ti.set_gdb_trigger(False)``
+- Show more detailed log (TI_TRACE): ``export TI_LOG_LEVEL=trace``
+- To specify which GPU to use for CUDA: ``export CUDA_VISIBLE_DEVICES=0``
+- To specify which Arch to use: ``export TI_ARCH=cuda``
+- To print intermediate IR generated: ``export TI_PRINT_IR=1``

docs/hello.rst

@@ -3,13 +3,11 @@ Hello, world!
 
 We introduce the Taichi programming language through a very basic `fractal` example.
 
-If you haven't done so, please install Taichi via ``pip``.
-Depending on your hardware and OS, please execute one of the following commands:
+If you haven't done so, please install Taichi via ``pip``:
 
 .. code-block:: bash
 
   # Python 3.6+ needed
-
   python3 -m pip install taichi
 
 Now you are ready to run the Taichi code below (``python3 fractal.py``) to compute a
@@ -30,7 +28,7 @@ Now you are ready to run the Taichi code below (``python3 fractal.py``) to compu
 
   @ti.func
   def complex_sqr(z):
-    return ti.Vector([z[0] * z[0] - z[1] * z[1], z[1] * z[0] * 2])
+    return ti.Vector([z[0] ** 2 - z[1] ** 2, z[1] * z[0] * 2])
 
   @ti.kernel
   def paint(t: ti.f32):
@@ -65,16 +63,35 @@ You can also reuse the package management system, Python IDEs, and existing Pyth
 Portability
 -----------------
 
-Taichi supports both CPUs and NVIDIA GPUs.
+Taichi code can run on CPUs or GPUs. Depends on your platform:
 
 .. code-block:: python
 
-  # Run on GPU
+  # Run on NVIDIA GPU, require CUDA installed
   ti.init(arch=ti.cuda)
+  # Run on GPU, with OpenGL backend
+  ti.init(arch=ti.opengl)
+  # Run on GPU, with Metal backend, if you're on OS X
+  ti.init(arch=ti.metal)
   # Run on CPU (default)
   ti.init(arch=ti.x64)
 
-If the machine does not have CUDA support, Taichi will fall back to CPUs instead.
+.. note::
+    Supported archs on different platforms:
+
+    +----------+------+------+--------+-------+
+    | platform | CPU  | CUDA | OpenGL | Metal |
+    +==========+======+======+========+=======+
+    | Windows  | OK   | WIP  | WIP    | MISS  |
+    +----------+------+------+--------+-------+
+    | Linux    | OK   | OK   | OK     | MISS  |
+    +----------+------+------+--------+-------+
+    | Mac OS X | OK   | MISS | MISS   | OK    |
+    +----------+------+------+--------+-------+
+
+    (OK=supported, WIP=work in progress, MISS=not supported)
+
+    If the machine does not have CUDA support, Taichi will fall back to CPUs instead.
 
 .. note::
 
@@ -86,7 +103,7 @@ If the machine does not have CUDA support, Taichi will fall back to CPUs instead
   On other platforms Taichi will make use of its on-demand memory allocator to adaptively allocate memory.
 
 (Sparse) Tensors
--------
+----------------
 
 Taichi is a data-oriented programming language, where dense or spatially-sparse tensors are first-class citizens.
 See :ref:`sparse` for more details on sparse tensors.
@@ -110,7 +127,7 @@ You can also define Taichi **functions** with ``ti.func``, which can be called a
 
 .. warning::
 
-  Taichi kernels must be called in the Python-scope. I.e., **nested Taichi kernels are not supported**.
+  Taichi kernels must be called in the Python-scope. I.e., **nested kernels are not supported**.
   Nested functions are allowed. **Recursive functions are not supported for now**.
 
   Taichi functions can only be called in Taichi-scope.
@@ -171,9 +188,23 @@ In the fractal code above, ``for i, j in pixels`` loops over all the pixel coord
           for i in x:
             ...
 
-.. warning::
+.. note::
+    ``break`` is not supported in **outermost (parallelized)** loops:
 
-    Struct-for's must be at the outer-most scope of kernels.
+    .. code-block:: python
+
+      @ti.kernel
+      def foo():
+        for i in x:
+            ...
+            break # ERROR! You cannot break a parallelized loop!
+
+      @ti.kernel
+      def foo():
+        for i in x:
+            for j in y:
+                ...
+                break # OK
 
 
 Interacting with Python

docs/index.rst

@@ -6,58 +6,65 @@ The Taichi Programming Language
    :maxdepth: 3
 
    overview
+   hello
+
 
 .. toctree::
    :caption: Basic Concepts
    :maxdepth: 3
 
-   hello
-
    syntax
-
    type
-
-   linalg
-
    tensor_matrix
-
-   global_settings
-
+   atomic
    external
 
+
 .. toctree::
    :caption: Advanced Programming
    :maxdepth: 3
 
    meta
-
    data_layout
-
    sparse
-
    differentiable_programming
-
    odop
-
    compilation
-
    syntax_sugars
 
 
 .. toctree::
-   :caption: Miscellaneous
+   :caption: References
    :maxdepth: 3
 
-   utilities
+   scalar_tensor
+   vector
+   linalg
+
+
+.. toctree::
+   :caption: Contribution
+   :maxdepth: 1
+
    dev_install
    contributor_guide
    cpp_style
    internal
-   faq
+
+
+.. toctree::
+   :caption: Miscellaneous
+   :maxdepth: 3
+
+   utilities
+   global_settings
+   performance
    acknowledgments
+   faq
+
 
 .. toctree::
    :caption: Legacy
    :maxdepth: 3
 
-   installation
+   legacy_installation

docs/internal.rst

@@ -7,6 +7,7 @@ Vector type system
 
 Intermediate representation
 ---------------------------------------
+Use ``ti.init(print_ir=True)`` to print IR on the console.
 
 
 Code generation