gaspersic/triton
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

[GH-PAGES] Updated website

Browse Source
This commit is contained in:
Philippe Tillet
2021-03-29 11:59:18 -04:00
parent 64141f0fca
commit 5cefc81fce
30 changed files with 1437 additions and 55 deletions
Download Patch File Download Diff File Expand all files Collapse all files

BIN
_downloads/662999063954282841dc90b8945f85ce/tutorials_jupyter.zip
View File

Binary file not shown.

BIN
_downloads/763344228ae6bc253ed1a6cf586aa30d/tutorials_python.zip
View File

Binary file not shown.

6
_downloads/b51b68bc1c6b1a5e509f67800b6235af/03-matrix-multiplication.ipynb
View File

@@ -83,7 +83,7 @@
},
"outputs": [],
"source": [
"src = \"\"\"\n#define MAX_GROUP_SIZE 8\n\n__global__ void dot(TYPE* A, TYPE* B, TYPE* C, \n int M, int N, int K, \n int lda, int ldb, int ldc) {\n int pid = get_program_id(0);\n int grid_m = (M + MB - 1) / MB;\n int grid_n = (N + NB - 1) / NB;\n int width = MAX_GROUP_SIZE * grid_n;\n int group_id = pid / width;\n int group_size = min(grid_m - group_id * MAX_GROUP_SIZE, MAX_GROUP_SIZE);\n int pid_m = group_id * MAX_GROUP_SIZE + (pid % group_size);\n int pid_n = (pid % width) / (group_size);\n int rm[MB] = pid_m * MB + 0 ... MB;\n int rn[NB] = pid_n * NB + 0 ... NB;\n int rk[KB] = 0 ... KB;\n TYPE *pa[MB, KB] = A + (rk [newaxis, :] * 1 + rm[:, newaxis] * lda);\n TYPE *pb[KB, NB] = B + (rk[:, newaxis] * ldb + rn [newaxis, :] * 1);\n float acc[MB, NB] = 0;\n for (int k = K; k > 0; k -= KB) {\n acc += (*pa) @ (*pb);\n pa += KB * 1;\n pb += KB * ldb;\n }\n rm = pid_m * MB + 0 ... MB;\n rn = pid_n * NB + 0 ... NB;\n TYPE *pc[MB, NB] = C + (rm[:, newaxis] * ldc + rn[newaxis, :]);\n *? (rm[:, newaxis] < M && rn [newaxis, :] < N) pc = acc;\n}\n\"\"\"\n\n\ndef make_kernel(device, dtype):\n key = (device, dtype)\n cache = make_kernel.cache\n if key not in cache:\n defines = {'TYPE': dtype}\n cache[key] = triton.kernel(src, device=device, defines=defines, autotune_vals=autotune_configs, autotune_key=autotune_key)\n return cache[key]\n\n\nmake_kernel.cache = dict()"
"src = \"\"\"\n#define MAX_GROUP_SIZE 8\n\n__global__ void dot(TYPE* A, TYPE* B, TYPE* C, \n int M, int N, int K, \n int lda, int ldb, int ldc) {\n int pid = get_program_id(0);\n int grid_m = (M + MB - 1) / MB;\n int grid_n = (N + NB - 1) / NB;\n int width = MAX_GROUP_SIZE * grid_n;\n int group_id = pid / width;\n int group_size = min(grid_m - group_id * MAX_GROUP_SIZE, MAX_GROUP_SIZE);\n int pid_m = group_id * MAX_GROUP_SIZE + (pid % group_size);\n int pid_n = (pid % width) / (group_size);\n int rm[MB] = pid_m * MB + 0 ... MB;\n int rn[NB] = pid_n * NB + 0 ... NB;\n int rk[KB] = 0 ... KB;\n TYPE *pa[MB, KB] = A + (rk [newaxis, :] * 1 + rm[:, newaxis] * lda);\n TYPE *pb[KB, NB] = B + (rk[:, newaxis] * ldb + rn [newaxis, :] * 1);\n float acc[MB, NB] = 0;\n for (int k = K; k > 0; k -= KB) {\n acc += (*pa) @ (*pb);\n pa += KB * 1;\n pb += KB * ldb;\n }\n rm = pid_m * MB + 0 ... MB;\n rn = pid_n * NB + 0 ... NB;\n TYPE *pc[MB, NB] = C + (rm[:, newaxis] * ldc + rn[newaxis, :]);\n *? (rm[:, newaxis] < M && rn [newaxis, :] < N) pc = acc;\n}\n\"\"\"\n\n\ndef make_kernel(device, dtype):\n key = (device, dtype)\n cache = make_kernel.cache\n if key not in cache:\n defines = {'TYPE': dtype}\n cache[key] = triton.kernel(\n src,\n device=device,\n defines=defines,\n autotune_configs=autotune_configs,\n autotune_key=autotune_key,\n )\n return cache[key]\n\n\nmake_kernel.cache = dict()"
]
},
{
@@ -126,7 +126,7 @@
"cell_type": "markdown",
"metadata": {},
"source": [
"## Benchmark\n\n### Installing The CUTLASS Bindings\n\nThe cuBLAS library (used by :code:`torch.matmul`) uses handwritten assembly-level optimizations that cannot be replicated using publicly available tools.\nFor this reason, we will instead compare the performance of our kernel against `CUTLASS <https://github.com/NVIDIA/cutlass/>`_ , a highly optimized CUDA library for matrix multiplication written by NVIDIA themselves._\nTo install CUTLASS, you need a recent version of cmake:\n\n .. code-block:: bash\n\n cd /path/to/cutlass/\n git clone https://github.com/NVIDIA/cutlass.git\n cd cutlass\n mkdir build\n cd build\n wget https://github.com/Kitware/CMake/releases/download/v3.19.4/cmake-3.19.4-Linux-x86_64.tar.gz\n tar xzvf *.tar.gz\n\nYou can then install CUTLASS as follows for V100\n\n .. code-block:: bash\n\n ./cmake-3.19.4-Linux-x86_64/bin/cmake ../ -DCUTLASS_NVCC_ARCHS_ENABLED=70 -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_f16_s884gemm_f16_*_align8\n make -j8 install\n\nOr as follows for A100:\n\n .. code-block:: bash\n\n ./cmake-3.19.4-Linux-x86_64/bin/cmake ../ -DCUTLASS_NVCC_ARCHS_ENABLED=80 -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_f16_s16816gemm_*align8\n make -j8 install\n\nWhere you can change CUTLASS_LIBRARY_KERNELS as you desire. Here, we are only interested in FP16 tensor core performance.\nTriton comes with some basic Python bindings for benchmarking CUTLASS. These will be compiled when the environment variables :code:`CUTLASS_INCLUDE_DIR` and :code:`CUTLASS_LIBRARY_DIR` are set during the installation process.\nTo re-install Triton with the updated CUTLASS bindings, run the following command:\n\n.. code-block:: bash\n\n export CUTLASS_INCLUDE_DIR=/tmp/cutlass/build/install/include/\n export CUTLASS_LIBRARY_DIR=/tmp/cutlass/build/install/lib/a\n pip uninstall -y triton\n pip install -e \"git+https://github.com/ptillet/triton.git#egg=triton&subdirectory=python\"\n\nWhich we can test as follows:\n\n"
"## Benchmark\n\n### Installing The CUTLASS Bindings\n\nThe cuBLAS library (used by :code:`torch.matmul`) uses handwritten assembly-level optimizations that cannot be replicated using publicly available tools.\nFor this reason, we will instead compare the performance of our kernel against `CUTLASS <https://github.com/NVIDIA/cutlass/>`_ , a highly optimized CUDA library for matrix multiplication written by NVIDIA themselves._\nTo install CUTLASS, you need a recent version of cmake:\n\n .. code-block:: bash\n\n cd /path/to/cutlass/\n git clone https://github.com/NVIDIA/cutlass.git\n cd cutlass\n mkdir build\n cd build\n wget https://github.com/Kitware/CMake/releases/download/v3.19.4/cmake-3.19.4-Linux-x86_64.tar.gz\n tar xzvf *.tar.gz\n\nYou can then install CUTLASS as follows for V100\n\n .. code-block:: bash\n\n ./cmake-3.19.4-Linux-x86_64/bin/cmake ../ -DCUTLASS_NVCC_ARCHS_ENABLED=70 -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_f16_s884gemm_f16_*_align8\n make -j8 install\n\nOr as follows for A100:\n\n .. code-block:: bash\n\n ./cmake-3.19.4-Linux-x86_64/bin/cmake ../ -DCUTLASS_NVCC_ARCHS_ENABLED=80 -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_f16_s16816gemm_*align8\n make -j8 install\n\nWhere you can change CUTLASS_LIBRARY_KERNELS as you desire. Here, we are only interested in FP16 tensor core performance.\nTriton comes with some basic Python bindings for benchmarking CUTLASS. These will be compiled when the environment variables :code:`CUTLASS_INCLUDE_DIR` and :code:`CUTLASS_LIBRARY_DIR` are set during the installation process.\nTo re-install Triton with the updated CUTLASS bindings, run the following command:\n\n.. code-block:: bash\n\n export CUTLASS_INCLUDE_DIR=/tmp/cutlass/build/install/include/\n export CUTLASS_LIBRARY_DIR=/tmp/cutlass/build/install/lib/\n pip uninstall -y triton\n pip install -e \"git+https://github.com/ptillet/triton.git#egg=triton&subdirectory=python\"\n\nWhich we can test as follows:\n\n"
]
},
{
@@ -155,7 +155,7 @@
},
"outputs": [],
"source": [
"@triton.testing.perf_report(\n triton.testing.Benchmark(\n x_names=['M', 'N', 'K'], # argument names to use as an x-axis for the plot\n x_vals=[256 * i for i in range(2, 33)], # different possible values for `x_name`\n y_name='provider', # argument name whose value corresponds to a different line in the plot\n y_vals=['torch', 'triton', 'cutlass'], # possible keys for `y_name`\n y_lines=[\"Torch\", \"Triton\", 'CUTLASS'], # label name for the lines\n ylabel=\"TFLOPS\", # label name for the y-axis\n plot_name=\"matmul-performance\", # name for the plot. Used also as a file name for saving the plot.\n args={}\n )\n)\ndef benchmark(M, N, K, provider):\n a = torch.randn((M, K), device='cuda', dtype=torch.float16)\n b = torch.randn((K, N), device='cuda', dtype=torch.float16)\n if provider == 'torch':\n ms, min_ms, max_ms = triton.testing.do_bench(lambda: torch.matmul(a, b))\n if provider == 'triton':\n ms, min_ms, max_ms = triton.testing.do_bench(lambda: dot(a, b))\n if provider == 'cutlass':\n ms, min_ms, max_ms = triton.testing.do_bench(lambda: triton.testing.cutlass_matmul(a, b))\n perf = lambda ms: 2 * M * N * K * 1e-12 / (ms * 1e-3)\n return perf(ms), perf(max_ms), perf(min_ms)\n\n\nbenchmark.run(show_plots=True)"
"@triton.testing.perf_report(\n triton.testing.Benchmark(\n x_names=['M', 'N', 'K'], # argument names to use as an x-axis for the plot\n x_vals=[256 * i for i in range(2, 33)], # different possible values for `x_name`\n y_name='provider', # argument name whose value corresponds to a different line in the plot\n y_vals=['cublas', 'triton', 'cutlass'], # possible keys for `y_name`\n y_lines=[\"cuBLAS\", \"Triton\", 'CUTLASS'], # label name for the lines\n ylabel=\"TFLOPS\", # label name for the y-axis\n plot_name=\"matmul-performance\", # name for the plot. Used also as a file name for saving the plot.\n args={}\n )\n)\ndef benchmark(M, N, K, provider):\n a = torch.randn((M, K), device='cuda', dtype=torch.float16)\n b = torch.randn((K, N), device='cuda', dtype=torch.float16)\n if provider == 'cublas':\n ms, min_ms, max_ms = triton.testing.do_bench(lambda: torch.matmul(a, b))\n if provider == 'triton':\n ms, min_ms, max_ms = triton.testing.do_bench(lambda: dot(a, b))\n if provider == 'cutlass':\n ms, min_ms, max_ms = triton.testing.do_bench(lambda: triton.testing.cutlass_matmul(a, b))\n perf = lambda ms: 2 * M * N * K * 1e-12 / (ms * 1e-3)\n return perf(ms), perf(max_ms), perf(min_ms)\n\n\nbenchmark.run(show_plots=True)"
]
},
{

16
_downloads/d5fee5b55a64e47f1b5724ec39adf171/03-matrix-multiplication.py
View File

@@ -229,7 +229,13 @@ def make_kernel(device, dtype):
cache = make_kernel.cache
if key not in cache:
defines = {'TYPE': dtype}
cache[key] = triton.kernel(src, device=device, defines=defines, autotune_vals=autotune_configs, autotune_key=autotune_key)
cache[key] = triton.kernel(
src,
device=device,
defines=defines,
autotune_configs=autotune_configs,
autotune_key=autotune_key,
)
return cache[key]
@@ -319,7 +325,7 @@ print(torch.allclose(c_0, c_1, rtol=1e-3, atol=1e-3))
# .. code-block:: bash
#
# export CUTLASS_INCLUDE_DIR=/tmp/cutlass/build/install/include/
# export CUTLASS_LIBRARY_DIR=/tmp/cutlass/build/install/lib/a
# export CUTLASS_LIBRARY_DIR=/tmp/cutlass/build/install/lib/
# pip uninstall -y triton
# pip install -e "git+https://github.com/ptillet/triton.git#egg=triton&subdirectory=python"
#
@@ -343,8 +349,8 @@ print(torch.allclose(c_0, c_2, rtol=1e-3, atol=1e-3))
x_names=['M', 'N', 'K'], # argument names to use as an x-axis for the plot
x_vals=[256 * i for i in range(2, 33)], # different possible values for `x_name`
y_name='provider', # argument name whose value corresponds to a different line in the plot
y_vals=['torch', 'triton', 'cutlass'], # possible keys for `y_name`
y_lines=["Torch", "Triton", 'CUTLASS'], # label name for the lines
y_vals=['cublas', 'triton', 'cutlass'], # possible keys for `y_name`
y_lines=["cuBLAS", "Triton", 'CUTLASS'], # label name for the lines
ylabel="TFLOPS", # label name for the y-axis
plot_name="matmul-performance", # name for the plot. Used also as a file name for saving the plot.
args={}
@@ -353,7 +359,7 @@ print(torch.allclose(c_0, c_2, rtol=1e-3, atol=1e-3))
def benchmark(M, N, K, provider):
a = torch.randn((M, K), device='cuda', dtype=torch.float16)
b = torch.randn((K, N), device='cuda', dtype=torch.float16)
if provider == 'torch':
if provider == 'cublas':
ms, min_ms, max_ms = triton.testing.do_bench(lambda: torch.matmul(a, b))
if provider == 'triton':
ms, min_ms, max_ms = triton.testing.do_bench(lambda: dot(a, b))

BIN
_images/cuda-parallel-matmul1.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 9.5 KiB

BIN
_images/halide-iteration1.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 12 KiB

BIN
_images/polyhedral-iteration1.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 59 KiB

BIN
_images/sphx_glr_01-vector-add_001.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 30 KiB

After

Width:  |  Height:  |  Size: 27 KiB

BIN
_images/sphx_glr_01-vector-add_thumb.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 19 KiB

After

Width:  |  Height:  |  Size: 17 KiB

BIN
_images/sphx_glr_02-fused-softmax_001.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 38 KiB

After

Width:  |  Height:  |  Size: 35 KiB

BIN
_images/sphx_glr_02-fused-softmax_thumb.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 24 KiB

After

Width:  |  Height:  |  Size: 22 KiB

BIN
_images/sphx_glr_03-matrix-multiplication_001.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 49 KiB

After

Width:  |  Height:  |  Size: 46 KiB

BIN
_images/sphx_glr_03-matrix-multiplication_thumb.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 31 KiB

After

Width:  |  Height:  |  Size: 29 KiB

BIN
_images/triton-parallel-matmul1.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 3.0 KiB

4
_sources/getting-started/tutorials/01-vector-add.rst.txt
View File

@@ -258,7 +258,7 @@ We can now run the decorated function above. Pass `show_plots=True` to see the p
.. image:: /getting-started/tutorials/images/sphx_glr_01-vector-add_001.png
:alt: vector-add-performance
:alt: 01 vector add
:class: sphx-glr-single-img
@@ -268,7 +268,7 @@ We can now run the decorated function above. Pass `show_plots=True` to see the p
.. rst-class:: sphx-glr-timing
**Total running time of the script:** ( 0 minutes 7.756 seconds)
**Total running time of the script:** ( 0 minutes 9.497 seconds)
.. _sphx_glr_download_getting-started_tutorials_01-vector-add.py:

4
_sources/getting-started/tutorials/02-fused-softmax.rst.txt
View File

@@ -295,7 +295,7 @@ We will then compare its performance against (1) :code:`torch.softmax` and (2) t
.. image:: /getting-started/tutorials/images/sphx_glr_02-fused-softmax_001.png
:alt: softmax-performance
:alt: 02 fused softmax
:class: sphx-glr-single-img
@@ -314,7 +314,7 @@ In the above plot, we can see that:
.. rst-class:: sphx-glr-timing
**Total running time of the script:** ( 0 minutes 19.933 seconds)
**Total running time of the script:** ( 0 minutes 25.654 seconds)
.. _sphx_glr_download_getting-started_tutorials_02-fused-softmax.py:

40
_sources/getting-started/tutorials/03-matrix-multiplication.rst.txt
View File

@@ -238,7 +238,7 @@ Here, we want to re-tune our kernel only when the shape of input matrices change
We can now create an auto-tuned kernel by passing the `autotune_configs` and `autotune_key` lists to the constructor of the :code:`triton.kernel` class.
.. GENERATED FROM PYTHON SOURCE LINES 193-238
.. GENERATED FROM PYTHON SOURCE LINES 193-244
.. code-block:: default
@@ -281,7 +281,13 @@ We can now create an auto-tuned kernel by passing the `autotune_configs` and `au
cache = make_kernel.cache
if key not in cache:
defines = {'TYPE': dtype}
cache[key] = triton.kernel(src, device=device, defines=defines, autotune_vals=autotune_configs, autotune_key=autotune_key)
cache[key] = triton.kernel(
src,
device=device,
defines=defines,
autotune_configs=autotune_configs,
autotune_key=autotune_key,
)
return cache[key]
@@ -294,7 +300,7 @@ We can now create an auto-tuned kernel by passing the `autotune_configs` and `au
.. GENERATED FROM PYTHON SOURCE LINES 239-244
.. GENERATED FROM PYTHON SOURCE LINES 245-250
Autograd Function
~~~~~~~~~~~~~~~~~~
@@ -302,7 +308,7 @@ Autograd Function
Now we are ready to expose our auto-tuned kernel as a `torch.autograd.Function`.
To do so, we just need to define a `forward` function that takes a two tensors as input and returns a tensor as output.
.. GENERATED FROM PYTHON SOURCE LINES 244-265
.. GENERATED FROM PYTHON SOURCE LINES 250-271
.. code-block:: default
@@ -334,7 +340,7 @@ To do so, we just need to define a `forward` function that takes a two tensors a
.. GENERATED FROM PYTHON SOURCE LINES 266-271
.. GENERATED FROM PYTHON SOURCE LINES 272-277
Unit Test
-----------
@@ -342,7 +348,7 @@ Unit Test
We can test our custom matrix multiplication operation against cuBLAS (i.e., :code:`torch.matmul`).
Note that we need to modify the :code`atol` and :code:`rtol` parameters of `torch.allclose` to account for the fact that we are comparing FP16 tensors.
.. GENERATED FROM PYTHON SOURCE LINES 271-280
.. GENERATED FROM PYTHON SOURCE LINES 277-286
.. code-block:: default
@@ -386,7 +392,7 @@ Note that we need to modify the :code`atol` and :code:`rtol` parameters of `torc
.. GENERATED FROM PYTHON SOURCE LINES 281-327
.. GENERATED FROM PYTHON SOURCE LINES 287-333
Benchmark
--------------
@@ -429,13 +435,13 @@ To re-install Triton with the updated CUTLASS bindings, run the following comman
.. code-block:: bash
export CUTLASS_INCLUDE_DIR=/tmp/cutlass/build/install/include/
export CUTLASS_LIBRARY_DIR=/tmp/cutlass/build/install/lib/a
export CUTLASS_LIBRARY_DIR=/tmp/cutlass/build/install/lib/
pip uninstall -y triton
pip install -e "git+https://github.com/ptillet/triton.git#egg=triton&subdirectory=python"
Which we can test as follows:
.. GENERATED FROM PYTHON SOURCE LINES 327-333
.. GENERATED FROM PYTHON SOURCE LINES 333-339
.. code-block:: default
@@ -468,7 +474,7 @@ Which we can test as follows:
.. GENERATED FROM PYTHON SOURCE LINES 334-339
.. GENERATED FROM PYTHON SOURCE LINES 340-345
Note that this wrapper for CUTLASS was written for benchmarking purposes and is probably not production-ready.
@@ -476,7 +482,7 @@ Square Matrix Performance
~~~~~~~~~~~~~~~~~~~~~~~~~~
We can now compare the performance of our kernel against CUTLASS. Here we focus on square matrices, but feel free to arrange the script as you wish to compare any other matrix shape.#
.. GENERATED FROM PYTHON SOURCE LINES 339-368
.. GENERATED FROM PYTHON SOURCE LINES 345-374
.. code-block:: default
@@ -487,8 +493,8 @@ We can now compare the performance of our kernel against CUTLASS. Here we focus
x_names=['M', 'N', 'K'], # argument names to use as an x-axis for the plot
x_vals=[256 * i for i in range(2, 33)], # different possible values for `x_name`
y_name='provider', # argument name whose value corresponds to a different line in the plot
y_vals=['torch', 'triton', 'cutlass'], # possible keys for `y_name`
y_lines=["Torch", "Triton", 'CUTLASS'], # label name for the lines
y_vals=['cublas', 'triton', 'cutlass'], # possible keys for `y_name`
y_lines=["cuBLAS", "Triton", 'CUTLASS'], # label name for the lines
ylabel="TFLOPS", # label name for the y-axis
plot_name="matmul-performance", # name for the plot. Used also as a file name for saving the plot.
args={}
@@ -497,7 +503,7 @@ We can now compare the performance of our kernel against CUTLASS. Here we focus
def benchmark(M, N, K, provider):
a = torch.randn((M, K), device='cuda', dtype=torch.float16)
b = torch.randn((K, N), device='cuda', dtype=torch.float16)
if provider == 'torch':
if provider == 'cublas':
ms, min_ms, max_ms = triton.testing.do_bench(lambda: torch.matmul(a, b))
if provider == 'triton':
ms, min_ms, max_ms = triton.testing.do_bench(lambda: dot(a, b))
@@ -513,21 +519,21 @@ We can now compare the performance of our kernel against CUTLASS. Here we focus
.. image:: /getting-started/tutorials/images/sphx_glr_03-matrix-multiplication_001.png
:alt: matmul-performance
:alt: 03 matrix multiplication
:class: sphx-glr-single-img
.. GENERATED FROM PYTHON SOURCE LINES 369-369
.. GENERATED FROM PYTHON SOURCE LINES 375-375
As we can see, the performance of our kernel is pretty good. It is in fact faster than CUTLASS, and therefore probably comparable to the absolute best CUDA code an expert could write.
.. rst-class:: sphx-glr-timing
**Total running time of the script:** ( 1 minutes 6.502 seconds)
**Total running time of the script:** ( 1 minutes 5.861 seconds)
.. _sphx_glr_download_getting-started_tutorials_03-matrix-multiplication.py:

8
_sources/getting-started/tutorials/sg_execution_times.rst.txt
View File

@@ -5,12 +5,12 @@
Computation times
=================
**01:34.190** total execution time for **getting-started_tutorials** files:
**00:25.654** total execution time for **getting-started_tutorials** files:
+---------------------------------------------------------------------------------------------------------+-----------+--------+
| :ref:`sphx_glr_getting-started_tutorials_03-matrix-multiplication.py` (``03-matrix-multiplication.py``) | 01:06.502 | 0.0 MB |
| :ref:`sphx_glr_getting-started_tutorials_02-fused-softmax.py` (``02-fused-softmax.py``) | 00:25.654 | 0.0 MB |
+---------------------------------------------------------------------------------------------------------+-----------+--------+
| :ref:`sphx_glr_getting-started_tutorials_02-fused-softmax.py` (``02-fused-softmax.py``) | 00:19.933 | 0.0 MB |
| :ref:`sphx_glr_getting-started_tutorials_01-vector-add.py` (``01-vector-add.py``) | 00:00.000 | 0.0 MB |
+---------------------------------------------------------------------------------------------------------+-----------+--------+
| :ref:`sphx_glr_getting-started_tutorials_01-vector-add.py` (``01-vector-add.py``) | 00:07.756 | 0.0 MB |
| :ref:`sphx_glr_getting-started_tutorials_03-matrix-multiplication.py` (``03-matrix-multiplication.py``) | 00:00.000 | 0.0 MB |
+---------------------------------------------------------------------------------------------------------+-----------+--------+

69
_sources/programming-guide/introduction.rst.txt Normal file
View File

@@ -0,0 +1,69 @@
==============
Introduction
==============
--------------
Motivations
--------------
Over the past decade, Deep Neural Networks (DNNs) have emerged as an important class of Machine Learning (ML) models, capable of achieving state-of-the-art performance across many domains ranging from natural language processing [1]_ to computer vision [2]_ to computational neuroscience [3]_. The strength of these models lies in their hierarchical structure, composed of a sequence of parametric (e.g., convolutional) and non-parametric (e.g., rectified linearity) *layers*. This pattern, though notoriously computationally expensive, also generates a large amount of highly parallelizable work particularly well suited for multi- and many- core processors.
As a consequence, Graphics Processing Units (GPUs) have become a cheap and accessible resource for exploring and/or deploying novel research ideas in the field. This trend has been accelerated by the release of several frameworks for General-Purpose GPU (GPGPU) computing, such as CUDA and OpenCL, which have made the development of high-performance programs easier. Yet, GPUs remain incredibly challenging to optimize for locality and parallelism, especially for computations that cannot be efficiently implemented using a combination of pre-existing optimized primitives. To make matters worse, GPU architectures are also rapidly evolving and specializing, as evidenced by the addition of tensor cores to NVIDIA (and more recently AMD) micro-architectures.
This tension between the computational opportunities offered by DNNs and the practical difficulty of GPU programming has created substantial academic and industrial interest for Domain-Specific Languages (DSLs) and compilers. Regrettably, these systems -- whether they be based on polyhedral machinery (*e.g.*, Tiramisu [4]_, Tensor Comprehensions [5]_) or scheduling languages (*e.g.*, Halide [6]_, TVM [7]_) -- remain less flexible and (for the same algorithm) markedly slower than the best handwritten compute kernels available in libraries like `cuBLAS <https://docs.nvidia.com/cuda/cublas/index.html>`_, `cuDNN <https://docs.nvidia.com/deeplearning/cudnn/api/index.html>`_ or `TensorRT <https://docs.nvidia.com/deeplearning/tensorrt/developer-guide/index.html>`_.
The main premise of this project is the following: programming paradigms based on blocked algorithms [8]_ can facilitate the construction of high-performance compute kernels for neural networks. We specifically revisit traditional "Single Program, Multiple Data" (SPMD [9]_) execution models for GPUs, and propose a variant in which programs -- rather than threads -- are blocked. For example, in the case of matrix multiplication, CUDA and Triton differ as follows:
.. table::
:widths: 50 50
+-----------------------------------------------------+-----------------------------------------------------+
| CUDA Programming Model | Triton Programming Model |
| | |
| (Scalar Program, Blocked Threads) | (Blocked Program, Scalar Threads) |
+=====================================================+=====================================================+
| | |
|.. code-block:: C |.. code-block:: C |
| | :force: |
| | |
| #pragma parallel | #pragma parallel |
| for(int m = 0; i < M; m++) | for(int m = 0; m < M; m += MB) |
| #pragma parallel | #pragma parallel |
| for(int n = 0; j < N; n++){ | for(int n = 0; n < N; n += NB){ |
| float acc = 0; | float acc[MB, NB] = 0; |
| for(int k = 0; k < K;k ++) | for(int k = 0; k < K; k += KB) |
| acc += A[i, k]* B[k, j]; | acc += A[m:m+MB, k:k+KB] |
| | @ B[k:k+KB, n:n+NB]; |
| C[i, j] = acc; | C[m:m+MB, n:n+NB] = acc; |
| } | } |
| | |
+-----------------------------------------------------+-----------------------------------------------------+
| |pic1| | |pic2| |
+-----------------------------------------------------+-----------------------------------------------------+
.. |pic1| image:: cuda-parallel-matmul.png
.. |pic2| image:: triton-parallel-matmul.png
A key benefit of this approach is that it leads to block-structured iteration spaces that offer programmers more flexibility than existing DSLs when implementing sparse operations, all while allowing compilers to aggressively optimize programs for data locality and parallelism.
--------------
Challenges
--------------
The main challenge posed by our proposed paradigm is that of work scheduling, i.e., how the work done by each program instance should be partitioned for efficient execution on modern GPUs. To address this issue, the Triton compiler makes heavy use of *block-level data-flow analysis*, a technique for scheduling iteration blocks statically based on the control- and data-flow structure of the target program. The resulting system actually works surprisingly well: our compiler manages to apply a broad range of interesting optimization automatically (e.g., automatic coalescing, thread swizzling, pre-fetching, automatic vectorization, tensor core-aware instruction selection, shared memory allocation/synchronization, asynchronous copy scheduling). Of course doing all this is not trivial; one of the purposes of this guide is to give you a sense of how it works.
--------------
References
--------------
.. [1] Sutskever et al., "Sequence to Sequence Learning with Neural Networks", NIPS 2014
.. [2] Redmon et al., "You Only Look Once: Unified, Real-Time Object Detection", CVPR 2016
.. [3] Lee et al., "Superhuman Accuracy on the SNEMI3D Connectomics Challenge", ArXiV 2017
.. [4] Baghdadi et al., "Tiramisu: A Polyhedral Compiler for Expressing Fast and Portable Code", CGO 2021
.. [5] Vasilache et al., "Tensor Comprehensions: Framework-Agnostic High-Performance Machine Learning Abstractions", ArXiV 2018
.. [6] Ragan-Kelley et al., "Halide: A Language and Compiler for Optimizing Parallelism, Locality, and Recomputation in Image Processing Pipelines", PLDI 2013
.. [7] Chen et al., "TVM: An Automated End-to-End Optimizing Compiler for Deep Learning", OSDI 2018
.. [8] Lam et al., "The Cache Performance and Optimizations of Blocked Algorithms", ASPLOS 1991
.. [9] Auguin et al., "Opsila: an advanced SIMD for numerical analysis and signal processing", EUROMICRO 1983

209
_sources/programming-guide/related-work.rst.txt Normal file
View File

@@ -0,0 +1,209 @@
==============
Related Work
==============
At first sight, Triton may seem like just yet another DSL for DNNs. The purpose of this section is to contextualize Triton and highlights its differences with the two leading approaches in this domain: polyhedral compilation and scheduling languages.
-----------------------
Polyhedral Compilation
-----------------------
Traditional compilers typically rely on intermediate representations, such as LLVM-IR [1]_, that encode control flow information using (un)conditional branches. This relatively low-level format makes it difficult to statically analyze the runtime behavior (e.g., cache misses) of input programs, and to automatically optimize loops accordingly through the use of tiling [2]_, fusion [3]_ and interchange [4]_. To solve this issue, polyhedral compilers [5]_ rely on program representations that have statically predictable control flow, thereby enabling aggressive compile-time program transformations for data locality and parallelism. Though this strategy has been adopted by many languages and compilers for DNNs such as Tiramisu [6]_, Tensor Comprehensions [7]_, Diesel [8]_ and the Affine dialect in MLIR [9]_, it also comes with a number of limitations that will be described later.
+++++++++++++++++++++++
Program Representation
+++++++++++++++++++++++
Polyhedral compilation is a vast area of research. In this section we only outline the most basic aspects of this topic, but readers interested in the solid mathematical foundations underneath may refer to the ample litterature on linear and integer programming.
.. table::
:widths: 50 50
+-----------------------------------------------------+-----------------------------------------------------+
| | |
|.. code-block:: C | |pic1| |
| | |
| for(int i = 0; i < 3; i++) | |
| for(int j = i; j < 5; j++) | |
| A[i][j] = 0; | |
+-----------------------------------------------------+-----------------------------------------------------+
.. |pic1| image:: polyhedral-iteration.png
:width: 300
Polyhedral compilers focus on a class of programs commonly known as **Static Control Parts** (SCoP), *i.e.*, maximal sets of consecutive statements in which conditionals and loop bounds are affine functions of surrounding loop indices and global invariant parameters. As shown above, programs in this format always lead to iteration domains that are bounded by affine inequalities, i.e., polyhedral. These polyhedra can also be defined algebraically; for the above example:
.. math::
\mathcal{P} = \{ i, j \in \mathbb{Z}^2
~|~
\begin{pmatrix}
1 & 0 \\
-1 & 0 \\
-1 & 1 \\
0 & -1 \\
\end{pmatrix}
\begin{pmatrix}
i \\
j
\end{pmatrix}
+
\begin{pmatrix}
0 \\
2 \\
0 \\
4
\end{pmatrix}
\geq
0
\}
Each point :math:`(i, j)` in :math:`\mathcal{P}` represents a *polyhedral statement*, that is a program statement which (1) does not induce control-flow side effects (e.g., :code:`for`, :code:`if`, :code:`break`) and (2) contains only affine functions of loop indices and global parameters in array accesses. To facilitate alias analysis, array accesses are also mathematically abstracted, using so-called *access function*. In other words, :code:`A[i][j]` is simply :code:`A[f(i,j)]` where the access function :math:`f` is defined by:
.. math::
f(i, j) = \begin{pmatrix}
1 & 0\\
0 & 1\\
\end{pmatrix}
\begin{pmatrix}
i\\
j
\end{pmatrix}
=
(i, j)
Note that the iteration domains of an SCoP does not specify the order in which its statements shall execute. In fact, this iteration domain may be traversed in many different possible legal orders, i.e. *schedules*. Formally, a schedule is defined as a p-dimensional affine transformation :math:`\Theta` of loop indices :math:`\mathbf{x}` and global invariant parameters :math:`\mathbf{g}`:
.. math::
\Theta_S(\mathbf{x}) = T_S \begin{pmatrix}
\vec{x}\\
\vec{g}\\
1
\end{pmatrix}
\qquad
T_S \in \mathbb{Z} ^{p \times (\text{dim}(\mathbf{x}) + \text{dim}(\mathbf{g}) + 1)}
Where :math:`\Theta_S(\mathbf{x})` is a p-dimensional vector representing the slowest to fastest growing indices (from left to right) when traversing the loop nest surrounding :math:`S`. For the code shown above, the original schedule defined by the loop nest in C can be retrieved by using:
.. math::
\Theta_S(\mathbf{x}) = \begin{pmatrix}
1 & 0 \\
0 & 1 \\
\end{pmatrix}
\begin{pmatrix}
i & j
\end{pmatrix}^T
=
\begin{pmatrix}
i & j
\end{pmatrix}^T
where :math:`i` and :math:`j` are respectively the slowest and fastest growing loop indices in the nest. If :math:`T_S` is a vector (resp. tensor), then :math:`\Theta_S` is a said to be one-dimensional (resp. multi-dimensional).
+++++++++++
Advantages
+++++++++++
Programs amenable to polyhedral compilation can be aggressively transformed and optimized. Most of these transformations actually boil down to the production of schedules and iteration domains that enable loop transformations promoting parallelism and spatial/temporal data locality (e.g., fusion, interchange, tiling, parallelization).
Polyhedral compilers can also automatically go through complex verification processes to ensure that the semantics of their input program is preserved throughout this optimization phase. Note that polyhedral optimizers are not incompatible with more standard optimization techniques. In fact, it is not uncommon for these systems to be implemented as a set of LLVM passes that can be run ahead of more traditional compilation techniques [10]_.
All in all, polyhedral machinery is extremely powerful, when applicable. It has been shown to support most common loop transformations, and has indeed achieved performance comparable to state-of-the-art GPU libraries for dense matrix multiplication [8]_. Additionally, it is also fully automatic and doesn't require any hint from programmers apart from source-code in a C-like format.
++++++++++++
Limitations
++++++++++++
Unfortunately, polyhedral compilers suffer from two major limitations that have prevented its adoption as a universal method for code generation in neural networks.
First, the set of possible program transformations $\Omega = \{ \Theta_S ~|~ S \in \text{program} \}$ is large, and grows with the number of statements in the program as well as with the size of their iteration domain. Verifying the legality of each transformation can also require the resolution of complex integer linear programs, making polyhedral compilation very computationally expensive. To make matters worse, hardware properties (e.g., cache size, number of SMs) and contextual characteristics (e.g., input tensor shapes) also have to be taken into account by this framework, leading to expensive auto-tuning procedures [11]_.
Second, the polyhedral framework is not very generally applicable; SCoPs are relatively common [12]_ but require loop bounds and array subscripts to be affine functions of loop indices, which typically only occurs in regular, dense computations. For this reason, this framework still has to be successfully applied to sparse -- or even structured-sparse -- neural networks, whose importance has been rapidly rising over the past few years.
On the other hand, blocked program representations advocated by this dissertation are less restricted in scope and can achieve close to peak performance using standard dataflow analysis.
-----------------------
Scheduling Languages
-----------------------
Separation of concerns \cite{dijkstra82} is a well-known design principle in computer science: programs should be decomposed into modular layers of abstraction that separate the semantics of their algorithms from the details of their implementation. Systems like Halide and TVM push this philosophy one step further, and enforce this separation at the grammatical level through the use of a **scheduling language**. The benefits of this methodology are particularly visible in the case of matrix multiplication, where, as one can see below, the definition of the algorithm (Line 1-7) is completely disjoint from its implementation (Line 8-16), meaning that both can be maintained, optimized and distributed independently.
.. code-block:: python
:linenos:
// algorithm
Var x("x"), y("y");
Func matmul("matmul");
RDom k(0, matrix_size);
RVar ki;
matmul(x, y) = 0.0f;
matmul(x, y) += A(k, y) * B(x, k);
// schedule
Var xi("xi"), xo("xo"), yo("yo"), yi("yo"), yii("yii"), xii("xii");
matmul.vectorize(x, 8);
matmul.update(0)
.split(x, x, xi, block_size).split(xi, xi, xii, 8)
.split(y, y, yi, block_size).split(yi, yi, yii, 4)
.split(k, k, ki, block_size)
.reorder(xii, yii, xi, ki, yi, k, x, y)
.parallel(y).vectorize(xii).unroll(xi).unroll(yii);
The resulting code may however not be completely portable, as schedules can sometimes rely on execution models (e.g., SPMD) or hardware intrinsics (e.g., matrix-multiply-accumulate) that are not widely available. This issue can be mitigated by auto-scheduling mechanisms [13]_.
+++++++++++
Advantages
+++++++++++
The main advantage of this approach is that it allows programmers to write an algorithm *only once*, and focus on performance optimization separately. It makes it possible to manually specify optimizations that a polyhedral compiler wouldn't be able to figure out automatically using static data-flow analysis.
Scheduling languages are, without a doubt, one of the most popular approaches for neural network code generation. The most popular system for this purpose is probably TVM, which provides good performance across a wide range of platforms as well as built-in automatic scheduling mechanisms.
++++++++++++
Limitations
++++++++++++
This ease-of-development comes at a cost. First of all, existing systems that follow this paradigm tend to be noticeably slower than Triton on modern hardware when applicable (e.g., V100/A100 tensor cores w/ equal tile sizes). I do believe that this is not a fundamental issue of scheduling languages -- in the sense that it could probably be solved with more efforts -- but it could mean that these systems are harder to engineer. More importantly, existing scheduling languages generate loops whose bounds and increments cannot depend on surrounding loop indice without at least imposing severe constraints on possible schedules -- if not breaking the system entirely. This is problematic for sparse com-putations, whose iteration spaces may be irregular.
.. table::
:widths: 50 50
+-----------------------------------------------------+-----------------------------------------------------+
| | |
|.. code-block:: C | |pic2| |
| | |
| for(int i = 0; i < 4; i++) | |
| for(int j = 0; j < 4; j++) | |
| float acc = 0; | |
| for(int k = 0; k < K[i]; k++) | |
| acc += A[i][col[i,k]]*B[k][j] | |
| C[i][j] = acc; | |
+-----------------------------------------------------+-----------------------------------------------------+
.. |pic2| image:: halide-iteration.png
:width: 300
On the other hand, the block-based program representation that we advocate for through this work allows for block-structured iteration spaces and allows programmers to manually handle load-balancing as they wish.
--------------
References
--------------
.. [1] Lattner et al., "LLVM: a compilation framework for lifelong program analysis transformation"
.. [2] Wolfe, "More Iteration Space Tiling", SC 1989
.. [3] Darte, "On the Complexity of Loop Fusion", PACT 1999
.. [4] Allen et al., "Automatic Loop Interchange", SIGPLAN Notices 1984
.. [5] Ancourt et al., "Scanning Polyhedra with DO Loops", PPoPP 1991
.. [6] Baghdadi et al., "Tiramisu: A Polyhedral Compiler for Expressing Fast and Portable Code", CGO 2021
.. [7] Vasilache et al., "Tensor Comprehensions: Framework-Agnostic High-Performance Machine Learning Abstractions", ArXiV 2018
.. [8] Elango et al. "Diesel: DSL for Linear Algebra and Neural Net Computations on GPUs", MAPL 2018
.. [9] Lattner et al., "MLIR Primer: A Compiler Infrastructure for the End of Moore’s Law", Arxiv 2019
.. [10] Grosser et al., "Polly - Performing Polyhedral Optimizations on a Low-Level Intermediate Representation", Parallel Processing Letters 2012
.. [11] Sato et al., "An Autotuning Framework for Scalable Execution of Tiled Code via Iterative Polyhedral Compilation", TACO 2019
.. [12] Girbal et al., "Semi-Automatic Composition of Loop Transformations for Deep Parallelism and Memory Hierarchies", International Journal of Parallel Programming 2006
.. [13] Mullapudi et al., "Automatically scheduling halide image processing pipelines", TOG 2016

83
_sources/programming-guide/triton-c.rst.txt Normal file
View File

@@ -0,0 +1,83 @@
=======================
The Triton-C Language
=======================
In the introduction, we stressed the importance of blocked algorithms and described their core principles in pseudo-code. To facilitate their implementation on modern GPU hardware, we present Triton-C, a single-threaded imperative kernel language in which block variables are first-class citizen. This language may be used either directly by developers familiar with C, or as an intermediate language for existing (and future) transcompilers. In this chapter, we describe its differences with C, its Numpy-like semantics and its "Single-Program, Multiple-Data" (SPMD) programming model.
-------------------
Differences with C
-------------------
The syntax of Triton-C is based on that of ANSI C, but was modified and extended to accomodate the semantics and programming model described in the next two subsections. These changes fall into the following categories:
+++++++++++
Extensions
+++++++++++
**Variable declarations**: Triton adds special-purpose syntax for multi-dimensional array declarations (e.g., :code:`int block[16, 16]`), which purposely differs from that of nested arrays (i.e., arrays of pointers) found in ANSI C (e.g., :code:`int block[16][16]`). Block dimensions must be constant but can also be made parametric with the use of pre-processor macros. One-dimensional blocks of integers may be initialized using ellipses (e.g., :code:`int range[16] = 0 ... 16`).
**Primitive types**: Triton-C supports the following primitive data-types: :code:`bool`, :code:`uint8`, :code:`uint16`, :code:`uint32`, :code:`uint64`, :code:`int8`, :code:`int16`, :code:`int32`, :code:`int64`, :code:`half`, :code:`float`, :code:`double`.
**Operators and built-in function**: The usual C operators were extended to support element-wise array operations (:code:`+`, :code:`-`, :code:`&&`, :code:`*`, etc.) and complex array operations(:code:`@` for matrix multiplication). Additionally, some built-in functions were added for concurrency (:code:`get_program_id`, :code:`atomic_add`).
**Slicing and broadcasting**: Multi-dimensional blocks can be broadcast along any particular dimension using numpy-like slicing syntax (e.g., :code:`int array[8, 8] = range[:, newaxis]` for stacking columns). Note that, as of now, slicing blocks to retrieve sub-blocks (or scalars) is forbidden as it is incompatible with the automatic parallelization methods used by our JIT. Reductions can be achieved using a syntax similar to slicing (e.g., :code:`array[+]` for summing an array, or :code:`array[:, max]` for row-wise maximum). Currently supported reduction operators are :code:`+`, :code:`min`, :code:`max`.
**Masked pointer dereferencement**: Block-level operations in Triton-C are "atomic", in the sense that they execute either completely or not at all. Basic element-wise control-flow for block-level operations can nonetheless be achieved using ternary operators and the *masked pointer dereferencement* operator exemplified below:
.. code-block:: C
// create mask
bool mask[16, 16] = ...;
// conditional addition
float x[16, 16] = mask ? a + b : 0;
// conditional load
float y[16] 16] = mask ? *ptr : 0;
// conditional store
*?(mask)ptr = y;
\end{lstlisting}
+++++++++++++
Restrictions
+++++++++++++
The Triton project is still in its infancy. As such, there are quite a few features of ANSI C that are not supported:
**Non-kernel functions**: Right now, all function definitions must be kernels, i.e. be preceded with the :code:`__global__` attribute. We are aware that this is a severe limitations, and the reason why it exists is because our automatic parallelization engine would not be capable of handling array parameter arguments.
**Non-primitive types**: Non-primitive types defined with :code:`struct` and :code:`union` are currently not supported, again because it is unclear at this point how these constructs would hook into our block-level data-flow analysis passes.
**While loops**: We just haven't had time to implement those yet.
----------------
Semantics
----------------
The existence of built-in **blocked** types, variable and operations in Triton-C offers two main benefits. First, it simplifies the structure of blocked programs by hiding important details pertaining to concurrent programming such as memory coalescing, cache management and specialized tensor instrinsics. Second, it opens the door for compilers to perform these optimizations automatically. However, it also means that programs have some kind of *block-level semantics* that does not exist in C. Though some aspects of it (e.g., the :code:`@` operator) are pretty intuitive, one in particular might be puzzling to some GPU programmers: broadcasting semantics.
+++++++++++++++++++++++
Broadcasting Semantics
+++++++++++++++++++++++
Block variables in Triton are strongly typed, meaning that certain instructions statically require their operands to satisfy strict shape constraints. For example, a scalar may not be added to an array unless it is first appropriately broadcast. *Broadcasting semantics* (first introduced in `Numpy <https://numpy.org/doc/stable/user/basics.broadcasting.html>`_) provides two formal rules for performing these conversions automatically in the case of binary operators: (1) the shape of the lowest-dimension operand is left-padded with ones until both operands have the same dimensionality; and (2) the content of both operands is replicated as many times as needed until their shape is identical. An error is emitted if this cannot be done.
.. code-block:: C
int a[16], b[32, 16], c[16, 1];
// a is first reshaped to [1, 16]
// and then broadcast to [32, 16]
int x_1[32, 16] = a[newaxis, :] + b;
// Same as above but implicitly
int x_2[32, 16] = a + b;
// a is first reshaped to [1, 16]
// a is broadcast to [16, 16]
// c is broadcast to [16, 16]
int y[16, 16] = a + c;
------------------
Programming Model
------------------
As discussed in the `CUDA documentation <https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html>`_, The execution of CUDA code on GPUs is supported by an `SPMD <https://en.wikipedia.org/wiki/SPMD>`_ programming model in which each kernel instance is associated with an identifiable *thread-block*, itself decomposed into *warps* of 32 *threads*. The Triton programming model is similar, but each kernel is *single-threaded* -- though automatically parallelized -- and associated with a global :code:`program id` which varies from instance to instance. This approach leads to simpler kernels in which CUDA-like concurrency primitives (shared memory synchronization, inter-thread communication, etc.) do not exist. The global program ids associated with each kernel instance can be queried using the :code:`get_program_id(axis)` built-in function where :code:`0 <= axis <= 2`. This is, for example, useful to create e.g., blocks of pointers as shown in the tutorials.

4
getting-started/tutorials/01-vector-add.html
View File

@@ -362,8 +362,8 @@ for different problem sizes.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="n">benchmark</span><span class="o">.</span><span class="n">run</span><span class="p">(</span><span class="n">show_plots</span><span class="o">=</span><span class="kc">True</span><span class="p">)</span>
</pre></div>
</div>
<img alt="vector-add-performance" class="sphx-glr-single-img" src="../../_images/sphx_glr_01-vector-add_001.png" />
<p class="sphx-glr-timing"><strong>Total running time of the script:</strong> ( 0 minutes 7.756 seconds)</p>
<img alt="01 vector add" class="sphx-glr-single-img" src="../../_images/sphx_glr_01-vector-add_001.png" />
<p class="sphx-glr-timing"><strong>Total running time of the script:</strong> ( 0 minutes 9.497 seconds)</p>
<div class="sphx-glr-footer class sphx-glr-footer-example docutils container" id="sphx-glr-download-getting-started-tutorials-01-vector-add-py">
<div class="sphx-glr-download sphx-glr-download-python docutils container">
<p><a class="reference download internal" download="" href="../../_downloads/62d97d49a32414049819dd8bb8378080/01-vector-add.py"><code class="xref download docutils literal notranslate"><span class="pre">Download</span> <span class="pre">Python</span> <span class="pre">source</span> <span class="pre">code:</span> <span class="pre">01-vector-add.py</span></code></a></p>

4
getting-started/tutorials/02-fused-softmax.html
View File

@@ -395,7 +395,7 @@ We will then compare its performance against (1) <code class="code docutils lite
<span class="n">benchmark</span><span class="o">.</span><span class="n">run</span><span class="p">(</span><span class="n">show_plots</span><span class="o">=</span><span class="kc">True</span><span class="p">)</span>
</pre></div>
</div>
<img alt="softmax-performance" class="sphx-glr-single-img" src="../../_images/sphx_glr_02-fused-softmax_001.png" />
<img alt="02 fused softmax" class="sphx-glr-single-img" src="../../_images/sphx_glr_02-fused-softmax_001.png" />
<p>In the above plot, we can see that:</p>
<blockquote>
<div><ul class="simple">
@@ -405,7 +405,7 @@ This means that – when temporary data is too large to fit entirely in the GPU
Note that our Triton kernel is not only faster than PyTorch’s CUDA kernel, it is also <strong>easier to read, understand and maintain</strong>.</p></li>
</ul>
</div></blockquote>
<p class="sphx-glr-timing"><strong>Total running time of the script:</strong> ( 0 minutes 19.933 seconds)</p>
<p class="sphx-glr-timing"><strong>Total running time of the script:</strong> ( 0 minutes 25.654 seconds)</p>
<div class="sphx-glr-footer class sphx-glr-footer-example docutils container" id="sphx-glr-download-getting-started-tutorials-02-fused-softmax-py">
<div class="sphx-glr-download sphx-glr-download-python docutils container">
<p><a class="reference download internal" download="" href="../../_downloads/d91442ac2982c4e0cc3ab0f43534afbc/02-fused-softmax.py"><code class="xref download docutils literal notranslate"><span class="pre">Download</span> <span class="pre">Python</span> <span class="pre">source</span> <span class="pre">code:</span> <span class="pre">02-fused-softmax.py</span></code></a></p>

20
getting-started/tutorials/03-matrix-multiplication.html
View File

@@ -411,7 +411,13 @@ Here, we want to re-tune our kernel only when the shape of input matrices change
<span class="n">cache</span> <span class="o">=</span> <span class="n">make_kernel</span><span class="o">.</span><span class="n">cache</span>
<span class="k">if</span> <span class="n">key</span> <span class="ow">not</span> <span class="ow">in</span> <span class="n">cache</span><span class="p">:</span>
<span class="n">defines</span> <span class="o">=</span> <span class="p">{</span><span class="s1">&#39;TYPE&#39;</span><span class="p">:</span> <span class="n">dtype</span><span class="p">}</span>
<span class="n">cache</span><span class="p">[</span><span class="n">key</span><span class="p">]</span> <span class="o">=</span> <span class="n">triton</span><span class="o">.</span><span class="n">kernel</span><span class="p">(</span><span class="n">src</span><span class="p">,</span> <span class="n">device</span><span class="o">=</span><span class="n">device</span><span class="p">,</span> <span class="n">defines</span><span class="o">=</span><span class="n">defines</span><span class="p">,</span> <span class="n">autotune_vals</span><span class="o">=</span><span class="n">autotune_configs</span><span class="p">,</span> <span class="n">autotune_key</span><span class="o">=</span><span class="n">autotune_key</span><span class="p">)</span>
<span class="n">cache</span><span class="p">[</span><span class="n">key</span><span class="p">]</span> <span class="o">=</span> <span class="n">triton</span><span class="o">.</span><span class="n">kernel</span><span class="p">(</span>
<span class="n">src</span><span class="p">,</span>
<span class="n">device</span><span class="o">=</span><span class="n">device</span><span class="p">,</span>
<span class="n">defines</span><span class="o">=</span><span class="n">defines</span><span class="p">,</span>
<span class="n">autotune_configs</span><span class="o">=</span><span class="n">autotune_configs</span><span class="p">,</span>
<span class="n">autotune_key</span><span class="o">=</span><span class="n">autotune_key</span><span class="p">,</span>
<span class="p">)</span>
<span class="k">return</span> <span class="n">cache</span><span class="p">[</span><span class="n">key</span><span class="p">]</span>
@@ -515,7 +521,7 @@ make -j8 install
Triton comes with some basic Python bindings for benchmarking CUTLASS. These will be compiled when the environment variables <code class="code docutils literal notranslate"><span class="pre">CUTLASS_INCLUDE_DIR</span></code> and <code class="code docutils literal notranslate"><span class="pre">CUTLASS_LIBRARY_DIR</span></code> are set during the installation process.
To re-install Triton with the updated CUTLASS bindings, run the following command:</p>
<div class="highlight-bash notranslate"><div class="highlight"><pre><span></span><span class="nb">export</span> <span class="nv">CUTLASS_INCLUDE_DIR</span><span class="o">=</span>/tmp/cutlass/build/install/include/
<span class="nb">export</span> <span class="nv">CUTLASS_LIBRARY_DIR</span><span class="o">=</span>/tmp/cutlass/build/install/lib/a
<span class="nb">export</span> <span class="nv">CUTLASS_LIBRARY_DIR</span><span class="o">=</span>/tmp/cutlass/build/install/lib/
pip uninstall -y triton
pip install -e <span class="s2">&quot;git+https://github.com/ptillet/triton.git#egg=triton&amp;subdirectory=python&quot;</span>
</pre></div>
@@ -549,8 +555,8 @@ True
<span class="n">x_names</span><span class="o">=</span><span class="p">[</span><span class="s1">&#39;M&#39;</span><span class="p">,</span> <span class="s1">&#39;N&#39;</span><span class="p">,</span> <span class="s1">&#39;K&#39;</span><span class="p">],</span> <span class="c1"># argument names to use as an x-axis for the plot</span>
<span class="n">x_vals</span><span class="o">=</span><span class="p">[</span><span class="mi">256</span> <span class="o">*</span> <span class="n">i</span> <span class="k">for</span> <span class="n">i</span> <span class="ow">in</span> <span class="nb">range</span><span class="p">(</span><span class="mi">2</span><span class="p">,</span> <span class="mi">33</span><span class="p">)],</span> <span class="c1"># different possible values for `x_name`</span>
<span class="n">y_name</span><span class="o">=</span><span class="s1">&#39;provider&#39;</span><span class="p">,</span> <span class="c1"># argument name whose value corresponds to a different line in the plot</span>
<span class="n">y_vals</span><span class="o">=</span><span class="p">[</span><span class="s1">&#39;torch&#39;</span><span class="p">,</span> <span class="s1">&#39;triton&#39;</span><span class="p">,</span> <span class="s1">&#39;cutlass&#39;</span><span class="p">],</span> <span class="c1"># possible keys for `y_name`</span>
<span class="n">y_lines</span><span class="o">=</span><span class="p">[</span><span class="s2">&quot;Torch&quot;</span><span class="p">,</span> <span class="s2">&quot;Triton&quot;</span><span class="p">,</span> <span class="s1">&#39;CUTLASS&#39;</span><span class="p">],</span> <span class="c1"># label name for the lines</span>
<span class="n">y_vals</span><span class="o">=</span><span class="p">[</span><span class="s1">&#39;cublas&#39;</span><span class="p">,</span> <span class="s1">&#39;triton&#39;</span><span class="p">,</span> <span class="s1">&#39;cutlass&#39;</span><span class="p">],</span> <span class="c1"># possible keys for `y_name`</span>
<span class="n">y_lines</span><span class="o">=</span><span class="p">[</span><span class="s2">&quot;cuBLAS&quot;</span><span class="p">,</span> <span class="s2">&quot;Triton&quot;</span><span class="p">,</span> <span class="s1">&#39;CUTLASS&#39;</span><span class="p">],</span> <span class="c1"># label name for the lines</span>
<span class="n">ylabel</span><span class="o">=</span><span class="s2">&quot;TFLOPS&quot;</span><span class="p">,</span> <span class="c1"># label name for the y-axis</span>
<span class="n">plot_name</span><span class="o">=</span><span class="s2">&quot;matmul-performance&quot;</span><span class="p">,</span> <span class="c1"># name for the plot. Used also as a file name for saving the plot.</span>
<span class="n">args</span><span class="o">=</span><span class="p">{}</span>
@@ -559,7 +565,7 @@ True
<span class="k">def</span> <span class="nf">benchmark</span><span class="p">(</span><span class="n">M</span><span class="p">,</span> <span class="n">N</span><span class="p">,</span> <span class="n">K</span><span class="p">,</span> <span class="n">provider</span><span class="p">):</span>
<span class="n">a</span> <span class="o">=</span> <span class="n">torch</span><span class="o">.</span><span class="n">randn</span><span class="p">((</span><span class="n">M</span><span class="p">,</span> <span class="n">K</span><span class="p">),</span> <span class="n">device</span><span class="o">=</span><span class="s1">&#39;cuda&#39;</span><span class="p">,</span> <span class="n">dtype</span><span class="o">=</span><span class="n">torch</span><span class="o">.</span><span class="n">float16</span><span class="p">)</span>
<span class="n">b</span> <span class="o">=</span> <span class="n">torch</span><span class="o">.</span><span class="n">randn</span><span class="p">((</span><span class="n">K</span><span class="p">,</span> <span class="n">N</span><span class="p">),</span> <span class="n">device</span><span class="o">=</span><span class="s1">&#39;cuda&#39;</span><span class="p">,</span> <span class="n">dtype</span><span class="o">=</span><span class="n">torch</span><span class="o">.</span><span class="n">float16</span><span class="p">)</span>
<span class="k">if</span> <span class="n">provider</span> <span class="o">==</span> <span class="s1">&#39;torch&#39;</span><span class="p">:</span>
<span class="k">if</span> <span class="n">provider</span> <span class="o">==</span> <span class="s1">&#39;cublas&#39;</span><span class="p">:</span>
<span class="n">ms</span><span class="p">,</span> <span class="n">min_ms</span><span class="p">,</span> <span class="n">max_ms</span> <span class="o">=</span> <span class="n">triton</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">do_bench</span><span class="p">(</span><span class="k">lambda</span><span class="p">:</span> <span class="n">torch</span><span class="o">.</span><span class="n">matmul</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">))</span>
<span class="k">if</span> <span class="n">provider</span> <span class="o">==</span> <span class="s1">&#39;triton&#39;</span><span class="p">:</span>
<span class="n">ms</span><span class="p">,</span> <span class="n">min_ms</span><span class="p">,</span> <span class="n">max_ms</span> <span class="o">=</span> <span class="n">triton</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">do_bench</span><span class="p">(</span><span class="k">lambda</span><span class="p">:</span> <span class="n">dot</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">))</span>
@@ -572,9 +578,9 @@ True
<span class="n">benchmark</span><span class="o">.</span><span class="n">run</span><span class="p">(</span><span class="n">show_plots</span><span class="o">=</span><span class="kc">True</span><span class="p">)</span>
</pre></div>
</div>
<img alt="matmul-performance" class="sphx-glr-single-img" src="../../_images/sphx_glr_03-matrix-multiplication_001.png" />
<img alt="03 matrix multiplication" class="sphx-glr-single-img" src="../../_images/sphx_glr_03-matrix-multiplication_001.png" />
<p>As we can see, the performance of our kernel is pretty good. It is in fact faster than CUTLASS, and therefore probably comparable to the absolute best CUDA code an expert could write.</p>
<p class="sphx-glr-timing"><strong>Total running time of the script:</strong> ( 1 minutes 6.502 seconds)</p>
<p class="sphx-glr-timing"><strong>Total running time of the script:</strong> ( 1 minutes 5.861 seconds)</p>
<div class="sphx-glr-footer class sphx-glr-footer-example docutils container" id="sphx-glr-download-getting-started-tutorials-03-matrix-multiplication-py">
<div class="sphx-glr-download sphx-glr-download-python docutils container">
<p><a class="reference download internal" download="" href="../../_downloads/d5fee5b55a64e47f1b5724ec39adf171/03-matrix-multiplication.py"><code class="xref download docutils literal notranslate"><span class="pre">Download</span> <span class="pre">Python</span> <span class="pre">source</span> <span class="pre">code:</span> <span class="pre">03-matrix-multiplication.py</span></code></a></p>

20
getting-started/tutorials/sg_execution_times.html
View File

@@ -167,7 +167,7 @@
<div class="section" id="computation-times">
<span id="sphx-glr-getting-started-tutorials-sg-execution-times"></span><h1>Computation times<a class="headerlink" href="#computation-times" title="Permalink to this headline"></a></h1>
<p><strong>01:34.190</strong> total execution time for <strong>getting-started_tutorials</strong> files:</p>
<p><strong>00:25.654</strong> total execution time for <strong>getting-started_tutorials</strong> files:</p>
<table class="docutils align-default">
<colgroup>
<col style="width: 85%" />
@@ -175,16 +175,16 @@
<col style="width: 6%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><a class="reference internal" href="02-fused-softmax.html#sphx-glr-getting-started-tutorials-02-fused-softmax-py"><span class="std std-ref">Fused Softmax</span></a> (<code class="docutils literal notranslate"><span class="pre">02-fused-softmax.py</span></code>)</p></td>
<td><p>00:25.654</p></td>
<td><p>0.0 MB</p></td>
</tr>
<tr class="row-even"><td><p><a class="reference internal" href="01-vector-add.html#sphx-glr-getting-started-tutorials-01-vector-add-py"><span class="std std-ref">Vector Addition</span></a> (<code class="docutils literal notranslate"><span class="pre">01-vector-add.py</span></code>)</p></td>
<td><p>00:00.000</p></td>
<td><p>0.0 MB</p></td>
</tr>
<tr class="row-odd"><td><p><a class="reference internal" href="03-matrix-multiplication.html#sphx-glr-getting-started-tutorials-03-matrix-multiplication-py"><span class="std std-ref">Matrix Multiplication</span></a> (<code class="docutils literal notranslate"><span class="pre">03-matrix-multiplication.py</span></code>)</p></td>
<td><p>01:06.502</p></td>
<td><p>0.0 MB</p></td>
</tr>
<tr class="row-even"><td><p><a class="reference internal" href="02-fused-softmax.html#sphx-glr-getting-started-tutorials-02-fused-softmax-py"><span class="std std-ref">Fused Softmax</span></a> (<code class="docutils literal notranslate"><span class="pre">02-fused-softmax.py</span></code>)</p></td>
<td><p>00:19.933</p></td>
<td><p>0.0 MB</p></td>
</tr>
<tr class="row-odd"><td><p><a class="reference internal" href="01-vector-add.html#sphx-glr-getting-started-tutorials-01-vector-add-py"><span class="std std-ref">Vector Addition</span></a> (<code class="docutils literal notranslate"><span class="pre">01-vector-add.py</span></code>)</p></td>
<td><p>00:07.756</p></td>
<td><p>00:00.000</p></td>
<td><p>0.0 MB</p></td>
</tr>
</tbody>

BIN
objects.inv
View File

Binary file not shown.

307
programming-guide/introduction.html Normal file
View File

@@ -0,0 +1,307 @@
<!DOCTYPE html>
<html class="writer-html5" lang="en" >
<head>
<meta charset="utf-8" />
<meta name="viewport" content="width=device-width, initial-scale=1.0" />
<title>Introduction &mdash; Triton documentation</title>
<link rel="stylesheet" href="../_static/css/theme.css" type="text/css" />
<link rel="stylesheet" href="../_static/pygments.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery-binder.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery-dataframe.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery-rendered-html.css" type="text/css" />
<link rel="stylesheet" href="../_static/css/custom.css" type="text/css" />
<!--[if lt IE 9]>
<script src="../_static/js/html5shiv.min.js"></script>
<![endif]-->
<script type="text/javascript" id="documentation_options" data-url_root="../" src="../_static/documentation_options.js"></script>
<script src="../_static/jquery.js"></script>
<script src="../_static/underscore.js"></script>
<script src="../_static/doctools.js"></script>
<script type="text/javascript" src="../_static/js/theme.js"></script>
<link rel="index" title="Index" href="../genindex.html" />
<link rel="search" title="Search" href="../search.html" />
</head>
<body class="wy-body-for-nav">
<div class="wy-grid-for-nav">
<nav data-toggle="wy-nav-shift" class="wy-nav-side">
<div class="wy-side-scroll">
<div class="wy-side-nav-search" >
<a href="../index.html" class="icon icon-home"> Triton
</a>
<div role="search">
<form id="rtd-search-form" class="wy-form" action="../search.html" method="get">
<input type="text" name="q" placeholder="Search docs" />
<input type="hidden" name="check_keywords" value="yes" />
<input type="hidden" name="area" value="default" />
</form>
</div>
</div>
<div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation">
<p class="caption"><span class="caption-text">Getting Started</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../getting-started/installation.html">Installation</a></li>
<li class="toctree-l1"><a class="reference internal" href="../getting-started/tutorials/index.html">Tutorials</a></li>
</ul>
<p class="caption"><span class="caption-text">Programming Guide</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="chapter-1/introduction.html">Introduction</a></li>
<li class="toctree-l1"><a class="reference internal" href="chapter-2/related-work.html">Related Work</a></li>
<li class="toctree-l1"><a class="reference internal" href="chapter-3/triton-c.html">The Triton-C Language</a></li>
<li class="toctree-l1"><a class="reference internal" href="chapter-4/triton-ir.html">The Triton-IR Intermediate Representation</a></li>
</ul>
</div>
</div>
</nav>
<section data-toggle="wy-nav-shift" class="wy-nav-content-wrap">
<nav class="wy-nav-top" aria-label="top navigation">
<i data-toggle="wy-nav-top" class="fa fa-bars"></i>
<a href="../index.html">Triton</a>
</nav>
<div class="wy-nav-content">
<div class="rst-content">
<div role="navigation" aria-label="breadcrumbs navigation">
<ul class="wy-breadcrumbs">
<li><a href="../index.html" class="icon icon-home"></a> &raquo;</li>
<li>Introduction</li>
<li class="wy-breadcrumbs-aside">
<a href="../_sources/programming-guide/introduction.rst.txt" rel="nofollow"> View page source</a>
</li>
</ul>
<hr/>
</div>
<div role="main" class="document" itemscope="itemscope" itemtype="http://schema.org/Article">
<div itemprop="articleBody">
<div class="section" id="introduction">
<h1>Introduction<a class="headerlink" href="#introduction" title="Permalink to this headline"></a></h1>
<div class="section" id="motivations">
<h2>Motivations<a class="headerlink" href="#motivations" title="Permalink to this headline"></a></h2>
<p>Over the past decade, Deep Neural Networks (DNNs) have emerged as an important class of Machine Learning (ML) models, capable of achieving state-of-the-art performance across many domains ranging from natural language processing <a class="footnote-reference brackets" href="#id10" id="id1">1</a> to computer vision <a class="footnote-reference brackets" href="#id11" id="id2">2</a> to computational neuroscience <a class="footnote-reference brackets" href="#id12" id="id3">3</a>. The strength of these models lies in their hierarchical structure, composed of a sequence of parametric (e.g., convolutional) and non-parametric (e.g., rectified linearity) <em>layers</em>. This pattern, though notoriously computationally expensive, also generates a large amount of highly parallelizable work particularly well suited for multi- and many- core processors.</p>
<p>As a consequence, Graphics Processing Units (GPUs) have become a cheap and accessible resource for exploring and/or deploying novel research ideas in the field. This trend has been accelerated by the release of several frameworks for General-Purpose GPU (GPGPU) computing, such as CUDA and OpenCL, which have made the development of high-performance programs easier. Yet, GPUs remain incredibly challenging to optimize for locality and parallelism, especially for computations that cannot be efficiently implemented using a combination of pre-existing optimized primitives. To make matters worse, GPU architectures are also rapidly evolving and specializing, as evidenced by the addition of tensor cores to NVIDIA (and more recently AMD) micro-architectures.</p>
<p>This tension between the computational opportunities offered by DNNs and the practical difficulty of GPU programming has created substantial academic and industrial interest for Domain-Specific Languages (DSLs) and compilers. Regrettably, these systems – whether they be based on polyhedral machinery (<em>e.g.</em>, Tiramisu <a class="footnote-reference brackets" href="#id13" id="id4">4</a>, Tensor Comprehensions <a class="footnote-reference brackets" href="#id14" id="id5">5</a>) or scheduling languages (<em>e.g.</em>, Halide <a class="footnote-reference brackets" href="#id15" id="id6">6</a>, TVM <a class="footnote-reference brackets" href="#id16" id="id7">7</a>) – remain less flexible and (for the same algorithm) markedly slower than the best handwritten compute kernels available in libraries like <a class="reference external" href="https://docs.nvidia.com/cuda/cublas/index.html">cuBLAS</a>, <a class="reference external" href="https://docs.nvidia.com/deeplearning/cudnn/api/index.html">cuDNN</a> or <a class="reference external" href="https://docs.nvidia.com/deeplearning/tensorrt/developer-guide/index.html">TensorRT</a>.</p>
<p>The main premise of this project is the following: programming paradigms based on blocked algorithms <a class="footnote-reference brackets" href="#id17" id="id8">8</a> can facilitate the construction of high-performance compute kernels for neural networks. We specifically revisit traditional “Single Program, Multiple Data” (SPMD <a class="footnote-reference brackets" href="#id18" id="id9">9</a>) execution models for GPUs, and propose a variant in which programs – rather than threads – are blocked. For example, in the case of matrix multiplication, CUDA and Triton differ as follows:</p>
<table class="colwidths-given docutils align-default">
<colgroup>
<col style="width: 50%" />
<col style="width: 50%" />
</colgroup>
<thead>
<tr class="row-odd"><th class="head"><p>CUDA Programming Model</p>
<p>(Scalar Program, Blocked Threads)</p>
</th>
<th class="head"><p>Triton Programming Model</p>
<p>(Blocked Program, Scalar Threads)</p>
</th>
</tr>
</thead>
<tbody>
<tr class="row-even"><td><div class="highlight-C notranslate"><div class="highlight"><pre><span></span><span class="cp">#pragma parallel</span>
<span class="k">for</span><span class="p">(</span><span class="kt">int</span> <span class="n">m</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span> <span class="n">i</span> <span class="o">&lt;</span> <span class="n">M</span><span class="p">;</span> <span class="n">m</span><span class="o">++</span><span class="p">)</span>
<span class="cp">#pragma parallel</span>
<span class="k">for</span><span class="p">(</span><span class="kt">int</span> <span class="n">n</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span> <span class="n">j</span> <span class="o">&lt;</span> <span class="n">N</span><span class="p">;</span> <span class="n">n</span><span class="o">++</span><span class="p">){</span>
<span class="kt">float</span> <span class="n">acc</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span>
<span class="k">for</span><span class="p">(</span><span class="kt">int</span> <span class="n">k</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span> <span class="n">k</span> <span class="o">&lt;</span> <span class="n">K</span><span class="p">;</span><span class="n">k</span> <span class="o">++</span><span class="p">)</span>
<span class="n">acc</span> <span class="o">+=</span> <span class="n">A</span><span class="p">[</span><span class="n">i</span><span class="p">,</span> <span class="n">k</span><span class="p">]</span><span class="o">*</span> <span class="n">B</span><span class="p">[</span><span class="n">k</span><span class="p">,</span> <span class="n">j</span><span class="p">];</span>
<span class="n">C</span><span class="p">[</span><span class="n">i</span><span class="p">,</span> <span class="n">j</span><span class="p">]</span> <span class="o">=</span> <span class="n">acc</span><span class="p">;</span>
<span class="p">}</span>
</pre></div>
</div>
</td>
<td><div class="highlight-C notranslate"><div class="highlight"><pre><span></span><span class="cp">#pragma parallel</span>
<span class="k">for</span><span class="p">(</span><span class="kt">int</span> <span class="n">m</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span> <span class="n">m</span> <span class="o">&lt;</span> <span class="n">M</span><span class="p">;</span> <span class="n">m</span> <span class="o">+=</span> <span class="n">MB</span><span class="p">)</span>
<span class="cp">#pragma parallel</span>
<span class="k">for</span><span class="p">(</span><span class="kt">int</span> <span class="n">n</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span> <span class="n">n</span> <span class="o">&lt;</span> <span class="n">N</span><span class="p">;</span> <span class="n">n</span> <span class="o">+=</span> <span class="n">NB</span><span class="p">){</span>
<span class="kt">float</span> <span class="n">acc</span><span class="p">[</span><span class="n">MB</span><span class="p">,</span> <span class="n">NB</span><span class="p">]</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span>
<span class="k">for</span><span class="p">(</span><span class="kt">int</span> <span class="n">k</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span> <span class="n">k</span> <span class="o">&lt;</span> <span class="n">K</span><span class="p">;</span> <span class="n">k</span> <span class="o">+=</span> <span class="n">KB</span><span class="p">)</span>
<span class="n">acc</span> <span class="o">+=</span> <span class="n">A</span><span class="p">[</span><span class="nl">m</span><span class="p">:</span><span class="n">m</span><span class="o">+</span><span class="n">MB</span><span class="p">,</span> <span class="nl">k</span><span class="p">:</span><span class="n">k</span><span class="o">+</span><span class="n">KB</span><span class="p">]</span>
<span class="err">@</span> <span class="n">B</span><span class="p">[</span><span class="nl">k</span><span class="p">:</span><span class="n">k</span><span class="o">+</span><span class="n">KB</span><span class="p">,</span> <span class="nl">n</span><span class="p">:</span><span class="n">n</span><span class="o">+</span><span class="n">NB</span><span class="p">];</span>
<span class="n">C</span><span class="p">[</span><span class="nl">m</span><span class="p">:</span><span class="n">m</span><span class="o">+</span><span class="n">MB</span><span class="p">,</span> <span class="nl">n</span><span class="p">:</span><span class="n">n</span><span class="o">+</span><span class="n">NB</span><span class="p">]</span> <span class="o">=</span> <span class="n">acc</span><span class="p">;</span>
<span class="p">}</span>
</pre></div>
</div>
</td>
</tr>
<tr class="row-odd"><td><p><img alt="pic1" src="../_images/cuda-parallel-matmul1.png" /></p></td>
<td><p><img alt="pic2" src="../_images/triton-parallel-matmul1.png" /></p></td>
</tr>
</tbody>
</table>
<p>A key benefit of this approach is that it leads to block-structured iteration spaces that offer programmers more flexibility than existing DSLs when implementing sparse operations, all while allowing compilers to aggressively optimize programs for data locality and parallelism.</p>
</div>
<div class="section" id="challenges">
<h2>Challenges<a class="headerlink" href="#challenges" title="Permalink to this headline"></a></h2>
<p>The main challenge posed by our proposed paradigm is that of work scheduling, i.e., how the work done by each program instance should be partitioned for efficient execution on modern GPUs. To address this issue, the Triton compiler makes heavy use of <em>block-level data-flow analysis</em>, a technique for scheduling iteration blocks statically based on the control- and data-flow structure of the target program. The resulting system actually works surprisingly well: our compiler manages to apply a broad range of interesting optimization automatically (e.g., automatic coalescing, thread swizzling, pre-fetching, automatic vectorization, tensor core-aware instruction selection, shared memory allocation/synchronization, asynchronous copy scheduling). Of course doing all this is not trivial; one of the purposes of this guide is to give you a sense of how it works.</p>
</div>
<div class="section" id="references">
<h2>References<a class="headerlink" href="#references" title="Permalink to this headline"></a></h2>
<dl class="footnote brackets">
<dt class="label" id="id10"><span class="brackets"><a class="fn-backref" href="#id1">1</a></span></dt>
<dd><p>Sutskever et al., “Sequence to Sequence Learning with Neural Networks”, NIPS 2014</p>
</dd>
<dt class="label" id="id11"><span class="brackets"><a class="fn-backref" href="#id2">2</a></span></dt>
<dd><p>Redmon et al., “You Only Look Once: Unified, Real-Time Object Detection”, CVPR 2016</p>
</dd>
<dt class="label" id="id12"><span class="brackets"><a class="fn-backref" href="#id3">3</a></span></dt>
<dd><p>Lee et al., “Superhuman Accuracy on the SNEMI3D Connectomics Challenge”, ArXiV 2017</p>
</dd>
<dt class="label" id="id13"><span class="brackets"><a class="fn-backref" href="#id4">4</a></span></dt>
<dd><p>Baghdadi et al., “Tiramisu: A Polyhedral Compiler for Expressing Fast and Portable Code”, CGO 2021</p>
</dd>
<dt class="label" id="id14"><span class="brackets"><a class="fn-backref" href="#id5">5</a></span></dt>
<dd><p>Vasilache et al., “Tensor Comprehensions: Framework-Agnostic High-Performance Machine Learning Abstractions”, ArXiV 2018</p>
</dd>
<dt class="label" id="id15"><span class="brackets"><a class="fn-backref" href="#id6">6</a></span></dt>
<dd><p>Ragan-Kelley et al., “Halide: A Language and Compiler for Optimizing Parallelism, Locality, and Recomputation in Image Processing Pipelines”, PLDI 2013</p>
</dd>
<dt class="label" id="id16"><span class="brackets"><a class="fn-backref" href="#id7">7</a></span></dt>
<dd><p>Chen et al., “TVM: An Automated End-to-End Optimizing Compiler for Deep Learning”, OSDI 2018</p>
</dd>
<dt class="label" id="id17"><span class="brackets"><a class="fn-backref" href="#id8">8</a></span></dt>
<dd><p>Lam et al., “The Cache Performance and Optimizations of Blocked Algorithms”, ASPLOS 1991</p>
</dd>
<dt class="label" id="id18"><span class="brackets"><a class="fn-backref" href="#id9">9</a></span></dt>
<dd><p>Auguin et al., “Opsila: an advanced SIMD for numerical analysis and signal processing”, EUROMICRO 1983</p>
</dd>
</dl>
</div>
</div>
</div>
</div>
<footer>
<hr/>
<div role="contentinfo">
<p>
&#169; Copyright 2020, Philippe Tillet.
</p>
</div>
Built with <a href="https://www.sphinx-doc.org/">Sphinx</a> using a
<a href="https://github.com/readthedocs/sphinx_rtd_theme">theme</a>
provided by <a href="https://readthedocs.org">Read the Docs</a>.
</footer>
</div>
</div>
</section>
</div>
<script type="text/javascript">
jQuery(function () {
SphinxRtdTheme.Navigation.enable(true);
});
</script>
</body>
</html>

425
programming-guide/related-work.html Normal file
View File

@@ -0,0 +1,425 @@
<!DOCTYPE html>
<html class="writer-html5" lang="en" >
<head>
<meta charset="utf-8" />
<meta name="viewport" content="width=device-width, initial-scale=1.0" />
<title>Related Work &mdash; Triton documentation</title>
<link rel="stylesheet" href="../_static/css/theme.css" type="text/css" />
<link rel="stylesheet" href="../_static/pygments.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery-binder.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery-dataframe.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery-rendered-html.css" type="text/css" />
<link rel="stylesheet" href="../_static/css/custom.css" type="text/css" />
<!--[if lt IE 9]>
<script src="../_static/js/html5shiv.min.js"></script>
<![endif]-->
<script type="text/javascript" id="documentation_options" data-url_root="../" src="../_static/documentation_options.js"></script>
<script src="../_static/jquery.js"></script>
<script src="../_static/underscore.js"></script>
<script src="../_static/doctools.js"></script>
<script async="async" src="https://cdnjs.cloudflare.com/ajax/libs/mathjax/2.7.7/latest.js?config=TeX-AMS-MML_HTMLorMML"></script>
<script type="text/javascript" src="../_static/js/theme.js"></script>
<link rel="index" title="Index" href="../genindex.html" />
<link rel="search" title="Search" href="../search.html" />
</head>
<body class="wy-body-for-nav">
<div class="wy-grid-for-nav">
<nav data-toggle="wy-nav-shift" class="wy-nav-side">
<div class="wy-side-scroll">
<div class="wy-side-nav-search" >
<a href="../index.html" class="icon icon-home"> Triton
</a>
<div role="search">
<form id="rtd-search-form" class="wy-form" action="../search.html" method="get">
<input type="text" name="q" placeholder="Search docs" />
<input type="hidden" name="check_keywords" value="yes" />
<input type="hidden" name="area" value="default" />
</form>
</div>
</div>
<div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation">
<p class="caption"><span class="caption-text">Getting Started</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../getting-started/installation.html">Installation</a></li>
<li class="toctree-l1"><a class="reference internal" href="../getting-started/tutorials/index.html">Tutorials</a></li>
</ul>
<p class="caption"><span class="caption-text">Programming Guide</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="chapter-1/introduction.html">Introduction</a></li>
<li class="toctree-l1"><a class="reference internal" href="chapter-2/related-work.html">Related Work</a></li>
<li class="toctree-l1"><a class="reference internal" href="chapter-3/triton-c.html">The Triton-C Language</a></li>
<li class="toctree-l1"><a class="reference internal" href="chapter-4/triton-ir.html">The Triton-IR Intermediate Representation</a></li>
</ul>
</div>
</div>
</nav>
<section data-toggle="wy-nav-shift" class="wy-nav-content-wrap">
<nav class="wy-nav-top" aria-label="top navigation">
<i data-toggle="wy-nav-top" class="fa fa-bars"></i>
<a href="../index.html">Triton</a>
</nav>
<div class="wy-nav-content">
<div class="rst-content">
<div role="navigation" aria-label="breadcrumbs navigation">
<ul class="wy-breadcrumbs">
<li><a href="../index.html" class="icon icon-home"></a> &raquo;</li>
<li>Related Work</li>
<li class="wy-breadcrumbs-aside">
<a href="../_sources/programming-guide/related-work.rst.txt" rel="nofollow"> View page source</a>
</li>
</ul>
<hr/>
</div>
<div role="main" class="document" itemscope="itemscope" itemtype="http://schema.org/Article">
<div itemprop="articleBody">
<div class="section" id="related-work">
<h1>Related Work<a class="headerlink" href="#related-work" title="Permalink to this headline"></a></h1>
<p>At first sight, Triton may seem like just yet another DSL for DNNs. The purpose of this section is to contextualize Triton and highlights its differences with the two leading approaches in this domain: polyhedral compilation and scheduling languages.</p>
<div class="section" id="polyhedral-compilation">
<h2>Polyhedral Compilation<a class="headerlink" href="#polyhedral-compilation" title="Permalink to this headline"></a></h2>
<p>Traditional compilers typically rely on intermediate representations, such as LLVM-IR <a class="footnote-reference brackets" href="#id17" id="id1">1</a>, that encode control flow information using (un)conditional branches. This relatively low-level format makes it difficult to statically analyze the runtime behavior (e.g., cache misses) of input programs, and to automatically optimize loops accordingly through the use of tiling <a class="footnote-reference brackets" href="#id18" id="id2">2</a>, fusion <a class="footnote-reference brackets" href="#id19" id="id3">3</a> and interchange <a class="footnote-reference brackets" href="#id20" id="id4">4</a>. To solve this issue, polyhedral compilers <a class="footnote-reference brackets" href="#id21" id="id5">5</a> rely on program representations that have statically predictable control flow, thereby enabling aggressive compile-time program transformations for data locality and parallelism. Though this strategy has been adopted by many languages and compilers for DNNs such as Tiramisu <a class="footnote-reference brackets" href="#id22" id="id6">6</a>, Tensor Comprehensions <a class="footnote-reference brackets" href="#id23" id="id7">7</a>, Diesel <a class="footnote-reference brackets" href="#id24" id="id8">8</a> and the Affine dialect in MLIR <a class="footnote-reference brackets" href="#id25" id="id9">9</a>, it also comes with a number of limitations that will be described later.</p>
<div class="section" id="program-representation">
<h3>Program Representation<a class="headerlink" href="#program-representation" title="Permalink to this headline"></a></h3>
<p>Polyhedral compilation is a vast area of research. In this section we only outline the most basic aspects of this topic, but readers interested in the solid mathematical foundations underneath may refer to the ample litterature on linear and integer programming.</p>
<table class="colwidths-given docutils align-default">
<colgroup>
<col style="width: 50%" />
<col style="width: 50%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><div class="highlight-C notranslate"><div class="highlight"><pre><span></span><span class="k">for</span><span class="p">(</span><span class="kt">int</span> <span class="n">i</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span> <span class="n">i</span> <span class="o">&lt;</span> <span class="mi">3</span><span class="p">;</span> <span class="n">i</span><span class="o">++</span><span class="p">)</span>
<span class="k">for</span><span class="p">(</span><span class="kt">int</span> <span class="n">j</span> <span class="o">=</span> <span class="n">i</span><span class="p">;</span> <span class="n">j</span> <span class="o">&lt;</span> <span class="mi">5</span><span class="p">;</span> <span class="n">j</span><span class="o">++</span><span class="p">)</span>
<span class="n">A</span><span class="p">[</span><span class="n">i</span><span class="p">][</span><span class="n">j</span><span class="p">]</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span>
</pre></div>
</div>
</td>
<td><p><a class="reference internal" href="../_images/polyhedral-iteration1.png"><img alt="pic1" src="../_images/polyhedral-iteration1.png" style="width: 300px;" /></a></p></td>
</tr>
</tbody>
</table>
<p>Polyhedral compilers focus on a class of programs commonly known as <strong>Static Control Parts</strong> (SCoP), <em>i.e.</em>, maximal sets of consecutive statements in which conditionals and loop bounds are affine functions of surrounding loop indices and global invariant parameters. As shown above, programs in this format always lead to iteration domains that are bounded by affine inequalities, i.e., polyhedral. These polyhedra can also be defined algebraically; for the above example:</p>
<div class="math notranslate nohighlight">
\[\begin{split}\mathcal{P} = \{ i, j \in \mathbb{Z}^2
~|~
\begin{pmatrix}
1 &amp; 0 \\
-1 &amp; 0 \\
-1 &amp; 1 \\
0 &amp; -1 \\
\end{pmatrix}
\begin{pmatrix}
i \\
j
\end{pmatrix}
+
\begin{pmatrix}
0 \\
2 \\
0 \\
4
\end{pmatrix}
\geq
0
\}\end{split}\]</div>
<p>Each point <span class="math notranslate nohighlight">\((i, j)\)</span> in <span class="math notranslate nohighlight">\(\mathcal{P}\)</span> represents a <em>polyhedral statement</em>, that is a program statement which (1) does not induce control-flow side effects (e.g., <code class="code docutils literal notranslate"><span class="pre">for</span></code>, <code class="code docutils literal notranslate"><span class="pre">if</span></code>, <code class="code docutils literal notranslate"><span class="pre">break</span></code>) and (2) contains only affine functions of loop indices and global parameters in array accesses. To facilitate alias analysis, array accesses are also mathematically abstracted, using so-called <em>access function</em>. In other words, <code class="code docutils literal notranslate"><span class="pre">A[i][j]</span></code> is simply <code class="code docutils literal notranslate"><span class="pre">A[f(i,j)]</span></code> where the access function <span class="math notranslate nohighlight">\(f\)</span> is defined by:</p>
<div class="math notranslate nohighlight">
\[\begin{split}f(i, j) = \begin{pmatrix}
1 &amp; 0\\
0 &amp; 1\\
\end{pmatrix}
\begin{pmatrix}
i\\
j
\end{pmatrix}
=
(i, j)\end{split}\]</div>
<p>Note that the iteration domains of an SCoP does not specify the order in which its statements shall execute. In fact, this iteration domain may be traversed in many different possible legal orders, i.e. <em>schedules</em>. Formally, a schedule is defined as a p-dimensional affine transformation <span class="math notranslate nohighlight">\(\Theta\)</span> of loop indices <span class="math notranslate nohighlight">\(\mathbf{x}\)</span> and global invariant parameters <span class="math notranslate nohighlight">\(\mathbf{g}\)</span>:</p>
<div class="math notranslate nohighlight">
\[\begin{split}\Theta_S(\mathbf{x}) = T_S \begin{pmatrix}
\vec{x}\\
\vec{g}\\
1
\end{pmatrix}
\qquad
T_S \in \mathbb{Z} ^{p \times (\text{dim}(\mathbf{x}) + \text{dim}(\mathbf{g}) + 1)}\end{split}\]</div>
<p>Where <span class="math notranslate nohighlight">\(\Theta_S(\mathbf{x})\)</span> is a p-dimensional vector representing the slowest to fastest growing indices (from left to right) when traversing the loop nest surrounding <span class="math notranslate nohighlight">\(S\)</span>. For the code shown above, the original schedule defined by the loop nest in C can be retrieved by using:</p>
<div class="math notranslate nohighlight">
\[\begin{split}\Theta_S(\mathbf{x}) = \begin{pmatrix}
1 &amp; 0 \\
0 &amp; 1 \\
\end{pmatrix}
\begin{pmatrix}
i &amp; j
\end{pmatrix}^T
=
\begin{pmatrix}
i &amp; j
\end{pmatrix}^T\end{split}\]</div>
<p>where <span class="math notranslate nohighlight">\(i\)</span> and <span class="math notranslate nohighlight">\(j\)</span> are respectively the slowest and fastest growing loop indices in the nest. If <span class="math notranslate nohighlight">\(T_S\)</span> is a vector (resp. tensor), then <span class="math notranslate nohighlight">\(\Theta_S\)</span> is a said to be one-dimensional (resp. multi-dimensional).</p>
</div>
<div class="section" id="advantages">
<h3>Advantages<a class="headerlink" href="#advantages" title="Permalink to this headline"></a></h3>
<p>Programs amenable to polyhedral compilation can be aggressively transformed and optimized. Most of these transformations actually boil down to the production of schedules and iteration domains that enable loop transformations promoting parallelism and spatial/temporal data locality (e.g., fusion, interchange, tiling, parallelization).</p>
<p>Polyhedral compilers can also automatically go through complex verification processes to ensure that the semantics of their input program is preserved throughout this optimization phase. Note that polyhedral optimizers are not incompatible with more standard optimization techniques. In fact, it is not uncommon for these systems to be implemented as a set of LLVM passes that can be run ahead of more traditional compilation techniques <a class="footnote-reference brackets" href="#id26" id="id10">10</a>.</p>
<p>All in all, polyhedral machinery is extremely powerful, when applicable. It has been shown to support most common loop transformations, and has indeed achieved performance comparable to state-of-the-art GPU libraries for dense matrix multiplication <a class="footnote-reference brackets" href="#id24" id="id11">8</a>. Additionally, it is also fully automatic and doesn’t require any hint from programmers apart from source-code in a C-like format.</p>
</div>
<div class="section" id="limitations">
<h3>Limitations<a class="headerlink" href="#limitations" title="Permalink to this headline"></a></h3>
<p>Unfortunately, polyhedral compilers suffer from two major limitations that have prevented its adoption as a universal method for code generation in neural networks.</p>
<p>First, the set of possible program transformations $Omega = { Theta_S ~|~ S in text{program} }$ is large, and grows with the number of statements in the program as well as with the size of their iteration domain. Verifying the legality of each transformation can also require the resolution of complex integer linear programs, making polyhedral compilation very computationally expensive. To make matters worse, hardware properties (e.g., cache size, number of SMs) and contextual characteristics (e.g., input tensor shapes) also have to be taken into account by this framework, leading to expensive auto-tuning procedures <a class="footnote-reference brackets" href="#id27" id="id12">11</a>.</p>
<p>Second, the polyhedral framework is not very generally applicable; SCoPs are relatively common <a class="footnote-reference brackets" href="#id28" id="id13">12</a> but require loop bounds and array subscripts to be affine functions of loop indices, which typically only occurs in regular, dense computations. For this reason, this framework still has to be successfully applied to sparse – or even structured-sparse – neural networks, whose importance has been rapidly rising over the past few years.</p>
<p>On the other hand, blocked program representations advocated by this dissertation are less restricted in scope and can achieve close to peak performance using standard dataflow analysis.</p>
</div>
</div>
<div class="section" id="scheduling-languages">
<h2>Scheduling Languages<a class="headerlink" href="#scheduling-languages" title="Permalink to this headline"></a></h2>
<p>Separation of concerns cite{dijkstra82} is a well-known design principle in computer science: programs should be decomposed into modular layers of abstraction that separate the semantics of their algorithms from the details of their implementation. Systems like Halide and TVM push this philosophy one step further, and enforce this separation at the grammatical level through the use of a <strong>scheduling language</strong>. The benefits of this methodology are particularly visible in the case of matrix multiplication, where, as one can see below, the definition of the algorithm (Line 1-7) is completely disjoint from its implementation (Line 8-16), meaning that both can be maintained, optimized and distributed independently.</p>
<div class="highlight-python notranslate"><table class="highlighttable"><tr><td class="linenos"><div class="linenodiv"><pre><span class="normal"> 1</span>
<span class="normal"> 2</span>
<span class="normal"> 3</span>
<span class="normal"> 4</span>
<span class="normal"> 5</span>
<span class="normal"> 6</span>
<span class="normal"> 7</span>
<span class="normal"> 8</span>
<span class="normal"> 9</span>
<span class="normal">10</span>
<span class="normal">11</span>
<span class="normal">12</span>
<span class="normal">13</span>
<span class="normal">14</span>
<span class="normal">15</span>
<span class="normal">16</span></pre></div></td><td class="code"><div class="highlight"><pre><span></span><span class="o">//</span> <span class="n">algorithm</span>
<span class="n">Var</span> <span class="n">x</span><span class="p">(</span><span class="s2">&quot;x&quot;</span><span class="p">),</span> <span class="n">y</span><span class="p">(</span><span class="s2">&quot;y&quot;</span><span class="p">);</span>
<span class="n">Func</span> <span class="n">matmul</span><span class="p">(</span><span class="s2">&quot;matmul&quot;</span><span class="p">);</span>
<span class="n">RDom</span> <span class="n">k</span><span class="p">(</span><span class="mi">0</span><span class="p">,</span> <span class="n">matrix_size</span><span class="p">);</span>
<span class="n">RVar</span> <span class="n">ki</span><span class="p">;</span>
<span class="n">matmul</span><span class="p">(</span><span class="n">x</span><span class="p">,</span> <span class="n">y</span><span class="p">)</span> <span class="o">=</span> <span class="mf">0.0</span><span class="n">f</span><span class="p">;</span>
<span class="n">matmul</span><span class="p">(</span><span class="n">x</span><span class="p">,</span> <span class="n">y</span><span class="p">)</span> <span class="o">+=</span> <span class="n">A</span><span class="p">(</span><span class="n">k</span><span class="p">,</span> <span class="n">y</span><span class="p">)</span> <span class="o">*</span> <span class="n">B</span><span class="p">(</span><span class="n">x</span><span class="p">,</span> <span class="n">k</span><span class="p">);</span>
<span class="o">//</span> <span class="n">schedule</span>
<span class="n">Var</span> <span class="n">xi</span><span class="p">(</span><span class="s2">&quot;xi&quot;</span><span class="p">),</span> <span class="n">xo</span><span class="p">(</span><span class="s2">&quot;xo&quot;</span><span class="p">),</span> <span class="n">yo</span><span class="p">(</span><span class="s2">&quot;yo&quot;</span><span class="p">),</span> <span class="n">yi</span><span class="p">(</span><span class="s2">&quot;yo&quot;</span><span class="p">),</span> <span class="n">yii</span><span class="p">(</span><span class="s2">&quot;yii&quot;</span><span class="p">),</span> <span class="n">xii</span><span class="p">(</span><span class="s2">&quot;xii&quot;</span><span class="p">);</span>
<span class="n">matmul</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">x</span><span class="p">,</span> <span class="mi">8</span><span class="p">);</span>
<span class="n">matmul</span><span class="o">.</span><span class="n">update</span><span class="p">(</span><span class="mi">0</span><span class="p">)</span>
<span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">x</span><span class="p">,</span> <span class="n">x</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">block_size</span><span class="p">)</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">xi</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">xii</span><span class="p">,</span> <span class="mi">8</span><span class="p">)</span>
<span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">y</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">yi</span><span class="p">,</span> <span class="n">block_size</span><span class="p">)</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">yi</span><span class="p">,</span> <span class="n">yi</span><span class="p">,</span> <span class="n">yii</span><span class="p">,</span> <span class="mi">4</span><span class="p">)</span>
<span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">k</span><span class="p">,</span> <span class="n">k</span><span class="p">,</span> <span class="n">ki</span><span class="p">,</span> <span class="n">block_size</span><span class="p">)</span>
<span class="o">.</span><span class="n">reorder</span><span class="p">(</span><span class="n">xii</span><span class="p">,</span> <span class="n">yii</span><span class="p">,</span> <span class="n">xi</span><span class="p">,</span> <span class="n">ki</span><span class="p">,</span> <span class="n">yi</span><span class="p">,</span> <span class="n">k</span><span class="p">,</span> <span class="n">x</span><span class="p">,</span> <span class="n">y</span><span class="p">)</span>
<span class="o">.</span><span class="n">parallel</span><span class="p">(</span><span class="n">y</span><span class="p">)</span><span class="o">.</span><span class="n">vectorize</span><span class="p">(</span><span class="n">xii</span><span class="p">)</span><span class="o">.</span><span class="n">unroll</span><span class="p">(</span><span class="n">xi</span><span class="p">)</span><span class="o">.</span><span class="n">unroll</span><span class="p">(</span><span class="n">yii</span><span class="p">);</span>
</pre></div>
</td></tr></table></div>
<p>The resulting code may however not be completely portable, as schedules can sometimes rely on execution models (e.g., SPMD) or hardware intrinsics (e.g., matrix-multiply-accumulate) that are not widely available. This issue can be mitigated by auto-scheduling mechanisms <a class="footnote-reference brackets" href="#id29" id="id14">13</a>.</p>
<div class="section" id="id15">
<h3>Advantages<a class="headerlink" href="#id15" title="Permalink to this headline"></a></h3>
<p>The main advantage of this approach is that it allows programmers to write an algorithm <em>only once</em>, and focus on performance optimization separately. It makes it possible to manually specify optimizations that a polyhedral compiler wouldn’t be able to figure out automatically using static data-flow analysis.</p>
<p>Scheduling languages are, without a doubt, one of the most popular approaches for neural network code generation. The most popular system for this purpose is probably TVM, which provides good performance across a wide range of platforms as well as built-in automatic scheduling mechanisms.</p>
</div>
<div class="section" id="id16">
<h3>Limitations<a class="headerlink" href="#id16" title="Permalink to this headline"></a></h3>
<p>This ease-of-development comes at a cost. First of all, existing systems that follow this paradigm tend to be noticeably slower than Triton on modern hardware when applicable (e.g., V100/A100 tensor cores w/ equal tile sizes). I do believe that this is not a fundamental issue of scheduling languages – in the sense that it could probably be solved with more efforts – but it could mean that these systems are harder to engineer. More importantly, existing scheduling languages generate loops whose bounds and increments cannot depend on surrounding loop indice without at least imposing severe constraints on possible schedules – if not breaking the system entirely. This is problematic for sparse com-putations, whose iteration spaces may be irregular.</p>
<table class="colwidths-given docutils align-default">
<colgroup>
<col style="width: 50%" />
<col style="width: 50%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><div class="highlight-C notranslate"><div class="highlight"><pre><span></span><span class="k">for</span><span class="p">(</span><span class="kt">int</span> <span class="n">i</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span> <span class="n">i</span> <span class="o">&lt;</span> <span class="mi">4</span><span class="p">;</span> <span class="n">i</span><span class="o">++</span><span class="p">)</span>
<span class="k">for</span><span class="p">(</span><span class="kt">int</span> <span class="n">j</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span> <span class="n">j</span> <span class="o">&lt;</span> <span class="mi">4</span><span class="p">;</span> <span class="n">j</span><span class="o">++</span><span class="p">)</span>
<span class="kt">float</span> <span class="n">acc</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span>
<span class="k">for</span><span class="p">(</span><span class="kt">int</span> <span class="n">k</span> <span class="o">=</span> <span class="mi">0</span><span class="p">;</span> <span class="n">k</span> <span class="o">&lt;</span> <span class="n">K</span><span class="p">[</span><span class="n">i</span><span class="p">];</span> <span class="n">k</span><span class="o">++</span><span class="p">)</span>
<span class="n">acc</span> <span class="o">+=</span> <span class="n">A</span><span class="p">[</span><span class="n">i</span><span class="p">][</span><span class="n">col</span><span class="p">[</span><span class="n">i</span><span class="p">,</span><span class="n">k</span><span class="p">]]</span><span class="o">*</span><span class="n">B</span><span class="p">[</span><span class="n">k</span><span class="p">][</span><span class="n">j</span><span class="p">]</span>
<span class="n">C</span><span class="p">[</span><span class="n">i</span><span class="p">][</span><span class="n">j</span><span class="p">]</span> <span class="o">=</span> <span class="n">acc</span><span class="p">;</span>
</pre></div>
</div>
</td>
<td><p><a class="reference internal" href="../_images/halide-iteration1.png"><img alt="pic2" src="../_images/halide-iteration1.png" style="width: 300px;" /></a></p></td>
</tr>
</tbody>
</table>
<p>On the other hand, the block-based program representation that we advocate for through this work allows for block-structured iteration spaces and allows programmers to manually handle load-balancing as they wish.</p>
</div>
</div>
<div class="section" id="references">
<h2>References<a class="headerlink" href="#references" title="Permalink to this headline"></a></h2>
<dl class="footnote brackets">
<dt class="label" id="id17"><span class="brackets"><a class="fn-backref" href="#id1">1</a></span></dt>
<dd><p>Lattner et al., “LLVM: a compilation framework for lifelong program analysis transformation”</p>
</dd>
<dt class="label" id="id18"><span class="brackets"><a class="fn-backref" href="#id2">2</a></span></dt>
<dd><p>Wolfe, “More Iteration Space Tiling”, SC 1989</p>
</dd>
<dt class="label" id="id19"><span class="brackets"><a class="fn-backref" href="#id3">3</a></span></dt>
<dd><p>Darte, “On the Complexity of Loop Fusion”, PACT 1999</p>
</dd>
<dt class="label" id="id20"><span class="brackets"><a class="fn-backref" href="#id4">4</a></span></dt>
<dd><p>Allen et al., “Automatic Loop Interchange”, SIGPLAN Notices 1984</p>
</dd>
<dt class="label" id="id21"><span class="brackets"><a class="fn-backref" href="#id5">5</a></span></dt>
<dd><p>Ancourt et al., “Scanning Polyhedra with DO Loops”, PPoPP 1991</p>
</dd>
<dt class="label" id="id22"><span class="brackets"><a class="fn-backref" href="#id6">6</a></span></dt>
<dd><p>Baghdadi et al., “Tiramisu: A Polyhedral Compiler for Expressing Fast and Portable Code”, CGO 2021</p>
</dd>
<dt class="label" id="id23"><span class="brackets"><a class="fn-backref" href="#id7">7</a></span></dt>
<dd><p>Vasilache et al., “Tensor Comprehensions: Framework-Agnostic High-Performance Machine Learning Abstractions”, ArXiV 2018</p>
</dd>
<dt class="label" id="id24"><span class="brackets">8</span><span class="fn-backref">(<a href="#id8">1</a>,<a href="#id11">2</a>)</span></dt>
<dd><p>Elango et al. “Diesel: DSL for Linear Algebra and Neural Net Computations on GPUs”, MAPL 2018</p>
</dd>
<dt class="label" id="id25"><span class="brackets"><a class="fn-backref" href="#id9">9</a></span></dt>
<dd><p>Lattner et al., “MLIR Primer: A Compiler Infrastructure for the End of Moore’s Law”, Arxiv 2019</p>
</dd>
<dt class="label" id="id26"><span class="brackets"><a class="fn-backref" href="#id10">10</a></span></dt>
<dd><p>Grosser et al., “Polly - Performing Polyhedral Optimizations on a Low-Level Intermediate Representation”, Parallel Processing Letters 2012</p>
</dd>
<dt class="label" id="id27"><span class="brackets"><a class="fn-backref" href="#id12">11</a></span></dt>
<dd><p>Sato et al., “An Autotuning Framework for Scalable Execution of Tiled Code via Iterative Polyhedral Compilation”, TACO 2019</p>
</dd>
<dt class="label" id="id28"><span class="brackets"><a class="fn-backref" href="#id13">12</a></span></dt>
<dd><p>Girbal et al., “Semi-Automatic Composition of Loop Transformations for Deep Parallelism and Memory Hierarchies”, International Journal of Parallel Programming 2006</p>
</dd>
<dt class="label" id="id29"><span class="brackets"><a class="fn-backref" href="#id14">13</a></span></dt>
<dd><p>Mullapudi et al., “Automatically scheduling halide image processing pipelines”, TOG 2016</p>
</dd>
</dl>
</div>
</div>
</div>
</div>
<footer>
<hr/>
<div role="contentinfo">
<p>
&#169; Copyright 2020, Philippe Tillet.
</p>
</div>
Built with <a href="https://www.sphinx-doc.org/">Sphinx</a> using a
<a href="https://github.com/readthedocs/sphinx_rtd_theme">theme</a>
provided by <a href="https://readthedocs.org">Read the Docs</a>.
</footer>
</div>
</div>
</section>
</div>
<script type="text/javascript">
jQuery(function () {
SphinxRtdTheme.Navigation.enable(true);
});
</script>
</body>
</html>

271
programming-guide/triton-c.html Normal file
View File

@@ -0,0 +1,271 @@
<!DOCTYPE html>
<html class="writer-html5" lang="en" >
<head>
<meta charset="utf-8" />
<meta name="viewport" content="width=device-width, initial-scale=1.0" />
<title>The Triton-C Language &mdash; Triton documentation</title>
<link rel="stylesheet" href="../_static/css/theme.css" type="text/css" />
<link rel="stylesheet" href="../_static/pygments.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery-binder.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery-dataframe.css" type="text/css" />
<link rel="stylesheet" href="../_static/gallery-rendered-html.css" type="text/css" />
<link rel="stylesheet" href="../_static/css/custom.css" type="text/css" />
<!--[if lt IE 9]>
<script src="../_static/js/html5shiv.min.js"></script>
<![endif]-->
<script type="text/javascript" id="documentation_options" data-url_root="../" src="../_static/documentation_options.js"></script>
<script src="../_static/jquery.js"></script>
<script src="../_static/underscore.js"></script>
<script src="../_static/doctools.js"></script>
<script type="text/javascript" src="../_static/js/theme.js"></script>
<link rel="index" title="Index" href="../genindex.html" />
<link rel="search" title="Search" href="../search.html" />
</head>
<body class="wy-body-for-nav">
<div class="wy-grid-for-nav">
<nav data-toggle="wy-nav-shift" class="wy-nav-side">
<div class="wy-side-scroll">
<div class="wy-side-nav-search" >
<a href="../index.html" class="icon icon-home"> Triton
</a>
<div role="search">
<form id="rtd-search-form" class="wy-form" action="../search.html" method="get">
<input type="text" name="q" placeholder="Search docs" />
<input type="hidden" name="check_keywords" value="yes" />
<input type="hidden" name="area" value="default" />
</form>
</div>
</div>
<div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation">
<p class="caption"><span class="caption-text">Getting Started</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../getting-started/installation.html">Installation</a></li>
<li class="toctree-l1"><a class="reference internal" href="../getting-started/tutorials/index.html">Tutorials</a></li>
</ul>
<p class="caption"><span class="caption-text">Programming Guide</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="chapter-1/introduction.html">Introduction</a></li>
<li class="toctree-l1"><a class="reference internal" href="chapter-2/related-work.html">Related Work</a></li>
<li class="toctree-l1"><a class="reference internal" href="chapter-3/triton-c.html">The Triton-C Language</a></li>
<li class="toctree-l1"><a class="reference internal" href="chapter-4/triton-ir.html">The Triton-IR Intermediate Representation</a></li>
</ul>
</div>
</div>
</nav>
<section data-toggle="wy-nav-shift" class="wy-nav-content-wrap">
<nav class="wy-nav-top" aria-label="top navigation">
<i data-toggle="wy-nav-top" class="fa fa-bars"></i>
<a href="../index.html">Triton</a>
</nav>
<div class="wy-nav-content">
<div class="rst-content">
<div role="navigation" aria-label="breadcrumbs navigation">
<ul class="wy-breadcrumbs">
<li><a href="../index.html" class="icon icon-home"></a> &raquo;</li>
<li>The Triton-C Language</li>
<li class="wy-breadcrumbs-aside">
<a href="../_sources/programming-guide/triton-c.rst.txt" rel="nofollow"> View page source</a>
</li>
</ul>
<hr/>
</div>
<div role="main" class="document" itemscope="itemscope" itemtype="http://schema.org/Article">
<div itemprop="articleBody">
<div class="section" id="the-triton-c-language">
<h1>The Triton-C Language<a class="headerlink" href="#the-triton-c-language" title="Permalink to this headline"></a></h1>
<p>In the introduction, we stressed the importance of blocked algorithms and described their core principles in pseudo-code. To facilitate their implementation on modern GPU hardware, we present Triton-C, a single-threaded imperative kernel language in which block variables are first-class citizen. This language may be used either directly by developers familiar with C, or as an intermediate language for existing (and future) transcompilers. In this chapter, we describe its differences with C, its Numpy-like semantics and its “Single-Program, Multiple-Data” (SPMD) programming model.</p>
<div class="section" id="differences-with-c">
<h2>Differences with C<a class="headerlink" href="#differences-with-c" title="Permalink to this headline"></a></h2>
<p>The syntax of Triton-C is based on that of ANSI C, but was modified and extended to accomodate the semantics and programming model described in the next two subsections. These changes fall into the following categories:</p>
<div class="section" id="extensions">
<h3>Extensions<a class="headerlink" href="#extensions" title="Permalink to this headline"></a></h3>
<p><strong>Variable declarations</strong>: Triton adds special-purpose syntax for multi-dimensional array declarations (e.g., <code class="code docutils literal notranslate"><span class="pre">int</span> <span class="pre">block[16,</span> <span class="pre">16]</span></code>), which purposely differs from that of nested arrays (i.e., arrays of pointers) found in ANSI C (e.g., <code class="code docutils literal notranslate"><span class="pre">int</span> <span class="pre">block[16][16]</span></code>). Block dimensions must be constant but can also be made parametric with the use of pre-processor macros. One-dimensional blocks of integers may be initialized using ellipses (e.g., <code class="code docutils literal notranslate"><span class="pre">int</span> <span class="pre">range[16]</span> <span class="pre">=</span> <span class="pre">0</span> <span class="pre">...</span> <span class="pre">16</span></code>).</p>
<p><strong>Primitive types</strong>: Triton-C supports the following primitive data-types: <code class="code docutils literal notranslate"><span class="pre">bool</span></code>, <code class="code docutils literal notranslate"><span class="pre">uint8</span></code>, <code class="code docutils literal notranslate"><span class="pre">uint16</span></code>, <code class="code docutils literal notranslate"><span class="pre">uint32</span></code>, <code class="code docutils literal notranslate"><span class="pre">uint64</span></code>, <code class="code docutils literal notranslate"><span class="pre">int8</span></code>, <code class="code docutils literal notranslate"><span class="pre">int16</span></code>, <code class="code docutils literal notranslate"><span class="pre">int32</span></code>, <code class="code docutils literal notranslate"><span class="pre">int64</span></code>, <code class="code docutils literal notranslate"><span class="pre">half</span></code>, <code class="code docutils literal notranslate"><span class="pre">float</span></code>, <code class="code docutils literal notranslate"><span class="pre">double</span></code>.</p>
<p><strong>Operators and built-in function</strong>: The usual C operators were extended to support element-wise array operations (<code class="code docutils literal notranslate"><span class="pre">+</span></code>, <code class="code docutils literal notranslate"><span class="pre">-</span></code>, <code class="code docutils literal notranslate"><span class="pre">&amp;&amp;</span></code>, <code class="code docutils literal notranslate"><span class="pre">*</span></code>, etc.) and complex array operations(<code class="code docutils literal notranslate"><span class="pre">&#64;</span></code> for matrix multiplication). Additionally, some built-in functions were added for concurrency (<code class="code docutils literal notranslate"><span class="pre">get_program_id</span></code>, <code class="code docutils literal notranslate"><span class="pre">atomic_add</span></code>).</p>
<p><strong>Slicing and broadcasting</strong>: Multi-dimensional blocks can be broadcast along any particular dimension using numpy-like slicing syntax (e.g., <code class="code docutils literal notranslate"><span class="pre">int</span> <span class="pre">array[8,</span> <span class="pre">8]</span> <span class="pre">=</span> <span class="pre">range[:,</span> <span class="pre">newaxis]</span></code> for stacking columns). Note that, as of now, slicing blocks to retrieve sub-blocks (or scalars) is forbidden as it is incompatible with the automatic parallelization methods used by our JIT. Reductions can be achieved using a syntax similar to slicing (e.g., <code class="code docutils literal notranslate"><span class="pre">array[+]</span></code> for summing an array, or <code class="code docutils literal notranslate"><span class="pre">array[:,</span> <span class="pre">max]</span></code> for row-wise maximum). Currently supported reduction operators are <code class="code docutils literal notranslate"><span class="pre">+</span></code>, <code class="code docutils literal notranslate"><span class="pre">min</span></code>, <code class="code docutils literal notranslate"><span class="pre">max</span></code>.</p>
<p><strong>Masked pointer dereferencement</strong>: Block-level operations in Triton-C are “atomic”, in the sense that they execute either completely or not at all. Basic element-wise control-flow for block-level operations can nonetheless be achieved using ternary operators and the <em>masked pointer dereferencement</em> operator exemplified below:</p>
<div class="highlight-C notranslate"><div class="highlight"><pre><span></span>// create mask
bool mask[16, 16] = ...;
// conditional addition
float x[16, 16] = mask ? a + b : 0;
// conditional load
float y[16] 16] = mask ? *ptr : 0;
// conditional store
*?(mask)ptr = y;
\end{lstlisting}
</pre></div>
</div>
</div>
<div class="section" id="restrictions">
<h3>Restrictions<a class="headerlink" href="#restrictions" title="Permalink to this headline"></a></h3>
<p>The Triton project is still in its infancy. As such, there are quite a few features of ANSI C that are not supported:</p>
<p><strong>Non-kernel functions</strong>: Right now, all function definitions must be kernels, i.e. be preceded with the <code class="code docutils literal notranslate"><span class="pre">__global__</span></code> attribute. We are aware that this is a severe limitations, and the reason why it exists is because our automatic parallelization engine would not be capable of handling array parameter arguments.</p>
<p><strong>Non-primitive types</strong>: Non-primitive types defined with <code class="code docutils literal notranslate"><span class="pre">struct</span></code> and <code class="code docutils literal notranslate"><span class="pre">union</span></code> are currently not supported, again because it is unclear at this point how these constructs would hook into our block-level data-flow analysis passes.</p>
<p><strong>While loops</strong>: We just haven’t had time to implement those yet.</p>
</div>
</div>
<div class="section" id="semantics">
<h2>Semantics<a class="headerlink" href="#semantics" title="Permalink to this headline"></a></h2>
<p>The existence of built-in <strong>blocked</strong> types, variable and operations in Triton-C offers two main benefits. First, it simplifies the structure of blocked programs by hiding important details pertaining to concurrent programming such as memory coalescing, cache management and specialized tensor instrinsics. Second, it opens the door for compilers to perform these optimizations automatically. However, it also means that programs have some kind of <em>block-level semantics</em> that does not exist in C. Though some aspects of it (e.g., the <code class="code docutils literal notranslate"><span class="pre">&#64;</span></code> operator) are pretty intuitive, one in particular might be puzzling to some GPU programmers: broadcasting semantics.</p>
<div class="section" id="broadcasting-semantics">
<h3>Broadcasting Semantics<a class="headerlink" href="#broadcasting-semantics" title="Permalink to this headline"></a></h3>
<p>Block variables in Triton are strongly typed, meaning that certain instructions statically require their operands to satisfy strict shape constraints. For example, a scalar may not be added to an array unless it is first appropriately broadcast. <em>Broadcasting semantics</em> (first introduced in <a class="reference external" href="https://numpy.org/doc/stable/user/basics.broadcasting.html">Numpy</a>) provides two formal rules for performing these conversions automatically in the case of binary operators: (1) the shape of the lowest-dimension operand is left-padded with ones until both operands have the same dimensionality; and (2) the content of both operands is replicated as many times as needed until their shape is identical. An error is emitted if this cannot be done.</p>
<div class="highlight-C notranslate"><div class="highlight"><pre><span></span><span class="kt">int</span> <span class="n">a</span><span class="p">[</span><span class="mi">16</span><span class="p">],</span> <span class="n">b</span><span class="p">[</span><span class="mi">32</span><span class="p">,</span> <span class="mi">16</span><span class="p">],</span> <span class="n">c</span><span class="p">[</span><span class="mi">16</span><span class="p">,</span> <span class="mi">1</span><span class="p">];</span>
<span class="c1">// a is first reshaped to [1, 16]</span>
<span class="c1">// and then broadcast to [32, 16]</span>
<span class="kt">int</span> <span class="n">x_1</span><span class="p">[</span><span class="mi">32</span><span class="p">,</span> <span class="mi">16</span><span class="p">]</span> <span class="o">=</span> <span class="n">a</span><span class="p">[</span><span class="n">newaxis</span><span class="p">,</span> <span class="o">:</span><span class="p">]</span> <span class="o">+</span> <span class="n">b</span><span class="p">;</span>
<span class="c1">// Same as above but implicitly</span>
<span class="kt">int</span> <span class="n">x_2</span><span class="p">[</span><span class="mi">32</span><span class="p">,</span> <span class="mi">16</span><span class="p">]</span> <span class="o">=</span> <span class="n">a</span> <span class="o">+</span> <span class="n">b</span><span class="p">;</span>
<span class="c1">// a is first reshaped to [1, 16]</span>
<span class="c1">// a is broadcast to [16, 16]</span>
<span class="c1">// c is broadcast to [16, 16]</span>
<span class="kt">int</span> <span class="n">y</span><span class="p">[</span><span class="mi">16</span><span class="p">,</span> <span class="mi">16</span><span class="p">]</span> <span class="o">=</span> <span class="n">a</span> <span class="o">+</span> <span class="n">c</span><span class="p">;</span>
</pre></div>
</div>
</div>
</div>
<div class="section" id="programming-model">
<h2>Programming Model<a class="headerlink" href="#programming-model" title="Permalink to this headline"></a></h2>
<p>As discussed in the <a class="reference external" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html">CUDA documentation</a>, The execution of CUDA code on GPUs is supported by an <a class="reference external" href="https://en.wikipedia.org/wiki/SPMD">SPMD</a> programming model in which each kernel instance is associated with an identifiable <em>thread-block</em>, itself decomposed into <em>warps</em> of 32 <em>threads</em>. The Triton programming model is similar, but each kernel is <em>single-threaded</em> – though automatically parallelized – and associated with a global <code class="code docutils literal notranslate"><span class="pre">program</span> <span class="pre">id</span></code> which varies from instance to instance. This approach leads to simpler kernels in which CUDA-like concurrency primitives (shared memory synchronization, inter-thread communication, etc.) do not exist. The global program ids associated with each kernel instance can be queried using the <code class="code docutils literal notranslate"><span class="pre">get_program_id(axis)</span></code> built-in function where <code class="code docutils literal notranslate"><span class="pre">0</span> <span class="pre">&lt;=</span> <span class="pre">axis</span> <span class="pre">&lt;=</span> <span class="pre">2</span></code>. This is, for example, useful to create e.g., blocks of pointers as shown in the tutorials.</p>
</div>
</div>
</div>
</div>
<footer>
<hr/>
<div role="contentinfo">
<p>
&#169; Copyright 2020, Philippe Tillet.
</p>
</div>
Built with <a href="https://www.sphinx-doc.org/">Sphinx</a> using a
<a href="https://github.com/readthedocs/sphinx_rtd_theme">theme</a>
provided by <a href="https://readthedocs.org">Read the Docs</a>.
</footer>
</div>
</div>
</section>
</div>
<script type="text/javascript">
jQuery(function () {
SphinxRtdTheme.Navigation.enable(true);
});
</script>
</body>
</html>

2
searchindex.js
View File

File diff suppressed because one or more lines are too long