[Doc] better documention

.travis.yml

@@ -20,6 +20,7 @@ matrix:
 skip_commits:
   files:
     - '*.md'
+    - '*.rst'
     - docs
     - benchmarks
     - examples

appveyor.yml

@@ -28,6 +28,7 @@ environment:
 skip_commits:
   files:
     - '*.md'
+    - '*.rst'
     - docs
     - benchmarks
     - examples

docs/atomic.rst

@@ -0,0 +1,74 @@
+.. _atomic:
+
+Atomic operations
+=================
+
+In Taichi, augmented assignments (e.g., ``x[i] += 1``) are automatically `atomic <https://en.wikipedia.org/wiki/Fetch-and-add>`_.
+
+
+.. warning::
+
+    When accumulating to global variables in parallel, make sure you use atomic operations. For example, to compute the sum of all elements in ``x``,
+    ::
+
+        @ti.kernel
+        def sum():
+            for i in x:
+                # Approach 1: OK
+                total[None] += x[i]
+
+                # Approach 2: OK
+                ti.atomic_add(total[None], x[i])
+
+                # Approach 3: Wrong result since the operation is not atomic.
+                total[None] = total[None] + x[i]
+
+
+.. note::
+    When atomic operations are applied to local values, the Taichi compiler will try to demote these operations into their non-atomic correspondence.
+
+Apart from augmented assignments, explicit atomic operations such as ``ti.atomic_add`` also do read-modify-write atomically.
+These operations additionally return the **old value** of the first argument. Below is the full list of explicit atomic operations:
+
+.. function:: ti.atomic_add(x, y)
+.. function:: ti.atomic_sub(x, y)
+
+    Atomically compute ``x + y``/``x - y`` and store the result to ``x``.
+
+    :return: The old value of ``x``.
+
+    For example,
+    ::
+
+        x = 3
+        y = 4
+        z = ti.atomic_add(x, y)
+        # now z = 7, y = 4, z = 3
+
+
+.. function:: ti.atomic_and(x, y)
+.. function:: ti.atomic_or(x, y)
+.. function:: ti.atomic_xor(x, y)
+
+    Atomically compute ``x & y`` (bitwise and), ``x | y`` (bitwise or), ``x ^ y`` (bitwise xor) and store the result to ``x``.
+
+    :return: The old value of ``x``.
+
+
+.. note::
+
+    Supported atomic operations on each backend:
+
+    +----------+-----------+-----------+---------+
+    | 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)

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/contributor_guide.rst

@@ -77,7 +77,7 @@ Using continuous integration
 ----------------------------
 
 - Continuous Integration (CI), will **build** and **test** your commits in a PR against in environments.
-- Currently, Taichi uses `"Travis CI" <https://travis-ci.org>`_(for OS X and Linux) and `"AppVeyor" <https://www.appveyor.com>`_(for Windows).
+- Currently, Taichi uses `Travis CI <https://travis-ci.org>`_ (for OS X and Linux) and `AppVeyor <https://www.appveyor.com>`_ (for Windows).
 - CI will be triggered everytime you push commits to an open PR.
 - You can prepend ``[skip ci]`` to your commit message to avoid triggering CI. e.g. ``[skip ci] This commit will not trigger CI``
 - A tick on the right of commit hash means CI passed, a cross means CI failed.

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 development
+---------------------------------------------
+
+- Clone the taichi repo, and build:
 
   .. code-block:: bash
 
+    git clone https://github.com/taichi-dev/taichi --depth=1 --branch=master
+    git submodule 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 the line below:
+    # 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 script 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 ``source ~/.bashrc`` to reload shell config.
+- Execute ``python3 -m taichi test`` to run all the tests. It may take up to 5 minutes to run all tests.
 - Check out ``examples`` for runnable examples. Run them with ``python3``.
 
 

docs/global_settings.rst

@@ -3,4 +3,11 @@ 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``
+- To not trigger GDB when crashes: ``export TI_GDB_TRIGGER=0``
+- To not use unified memory for CUDA: ``export TI_USE_UNIFIED_MEMORY=0``
+- To specify pre-allocated memory size for CUDA: ``export TI_DEVICE_MEMORY_GB=0.5``
+- 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``
+- To print verbosed details: ``export TI_VERBOSE=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. Initialize Taichi according to your hardware platform:
 
 .. code-block:: python
 
-  # Run on GPU
+  # Run on NVIDIA GPU, CUDA required
   ti.init(arch=ti.cuda)
+  # Run on GPU, with the OpenGL backend
+  ti.init(arch=ti.opengl)
+  # Run on GPU, with the Apple Metal backend, if you are 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 backends on different platforms:
+
+    +----------+------+------+--------+-------+
+    | platform | CPU  | CUDA | OpenGL | Metal |
+    +==========+======+======+========+=======+
+    | Windows  | OK   | OK   | WIP    | N/A   |
+    +----------+------+------+--------+-------+
+    | Linux    | OK   | OK   | OK     | N/A   |
+    +----------+------+------+--------+-------+
+    | Mac OS X | OK   | N/A  | N/A    | OK    |
+    +----------+------+------+--------+-------+
+
+    (OK: supported, WIP: work in progress, N/A: not available)
+
+    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.
@@ -152,6 +169,9 @@ In the fractal code above, ``for i, j in pixels`` loops over all the pixel coord
 
     Struct-for is the key to :ref:`sparse` in Taichi, as it will only loop over active elements in a sparse tensor. In dense tensors, all elements are active.
 
+.. note::
+    Struct-for's must be at the outer-most scope of kernels.
+
 .. note::
     It is the loop **at the outermost scope** that gets parallelized, not the outermost loop.
 
@@ -171,9 +191,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