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
2022-04-08 00:44:05 +00:00
parent 80b92a0d2d
commit 0c570c178d
173 changed files with 401 additions and 386 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -31,7 +31,7 @@ In this tutorial, you will write a simple vector addition using Triton and learn
Compute Kernel
--------------------------
.. GENERATED FROM PYTHON SOURCE LINES 14-50
.. GENERATED FROM PYTHON SOURCE LINES 14-53
.. code-block:: default
@@ -48,9 +48,11 @@ Compute Kernel
y_ptr, # *Pointer* to second input vector
output_ptr, # *Pointer* to output vector
n_elements, # Size of the vector
time_start_ptr, time_end_ptr,
BLOCK_SIZE: tl.constexpr, # Number of elements each program should process
# NOTE: `constexpr` so it can be used as a shape value
):
tl.atomic_min(time_start_ptr, tl.clock())
# There are multiple 'program's processing different data. We identify which program
# we are here
pid = tl.program_id(axis=0) # We use a 1D launch grid so axis is 0
@@ -69,6 +71,7 @@ Compute Kernel
output = x + y
# Write x + y back to DRAM
tl.store(output_ptr + offsets, output, mask=mask)
tl.atomic_max(time_end_ptr, tl.clock())
@@ -78,18 +81,20 @@ Compute Kernel
.. GENERATED FROM PYTHON SOURCE LINES 51-53
.. GENERATED FROM PYTHON SOURCE LINES 54-56
Let's also declare a helper function to (1) allocate the `z` tensor
and (2) enqueue the above kernel with appropriate grid/block sizes.
.. GENERATED FROM PYTHON SOURCE LINES 53-74
.. GENERATED FROM PYTHON SOURCE LINES 56-79
.. code-block:: default
def add(x: torch.Tensor, y: torch.Tensor):
time_start = torch.zeros(1, dtype=torch.int64, device='cuda')
time_end = torch.zeros(1, dtype=torch.int64, device='cuda')
# We need to preallocate the output
output = torch.empty_like(x)
assert x.is_cuda and y.is_cuda and output.is_cuda
@@ -102,7 +107,7 @@ and (2) enqueue the above kernel with appropriate grid/block sizes.
# - each torch.tensor object is implicitly converted into a pointer to its first element.
# - `triton.jit`'ed functions can be index with a launch grid to obtain a callable GPU kernel
# - don't forget to pass meta-parameters as keywords arguments
add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
add_kernel[grid](x, y, output, n_elements, time_start, time_end, BLOCK_SIZE=1024)
# We return a handle to z but, since `torch.cuda.synchronize()` hasn't been called, the kernel is still
# running asynchronously at this point.
return output
@@ -115,11 +120,11 @@ and (2) enqueue the above kernel with appropriate grid/block sizes.
.. GENERATED FROM PYTHON SOURCE LINES 75-76
.. GENERATED FROM PYTHON SOURCE LINES 80-81
We can now use the above function to compute the element-wise sum of two `torch.tensor` objects and test its correctness:
.. GENERATED FROM PYTHON SOURCE LINES 76-90
.. GENERATED FROM PYTHON SOURCE LINES 81-95
.. code-block:: default
@@ -154,11 +159,11 @@ We can now use the above function to compute the element-wise sum of two `torch.
.. GENERATED FROM PYTHON SOURCE LINES 91-92
.. GENERATED FROM PYTHON SOURCE LINES 96-97
Seems like we're good to go!
.. GENERATED FROM PYTHON SOURCE LINES 94-99
.. GENERATED FROM PYTHON SOURCE LINES 99-104
Benchmark
-----------
@@ -166,7 +171,7 @@ We can now benchmark our custom op on vectors of increasing sizes to get a sense
To make things easier, Triton has a set of built-in utilities that allow us to concisely plot the performance of your custom ops
for different problem sizes.
.. GENERATED FROM PYTHON SOURCE LINES 99-128
.. GENERATED FROM PYTHON SOURCE LINES 104-133
.. code-block:: default
@@ -206,12 +211,12 @@ for different problem sizes.
.. GENERATED FROM PYTHON SOURCE LINES 129-131
.. GENERATED FROM PYTHON SOURCE LINES 134-136
We can now run the decorated function above. Pass `print_data=True` to see the performance number, `show_plots=True` to plot them, and/or
`save_path='/path/to/results/' to save them to disk along with raw CSV data
.. GENERATED FROM PYTHON SOURCE LINES 131-132
.. GENERATED FROM PYTHON SOURCE LINES 136-137
.. code-block:: default
@@ -232,22 +237,22 @@ We can now run the decorated function above. Pass `print_data=True` to see the p
vector-add-performance:
size Triton Torch
0 4096.0 9.600000 9.600000
1 8192.0 19.200000 19.200000
2 16384.0 38.400001 38.400001
3 32768.0 76.800002 76.800002
4 65536.0 127.999995 127.999995
5 131072.0 219.428568 219.428568
6 262144.0 341.333321 341.333321
7 524288.0 472.615390 472.615390
8 1048576.0 614.400016 614.400016
9 2097152.0 722.823517 722.823517
10 4194304.0 780.190482 780.190482
11 8388608.0 812.429770 812.429770
12 16777216.0 833.084721 833.084721
13 33554432.0 842.004273 843.811163
14 67108864.0 847.448255 848.362445
15 134217728.0 849.737435 850.656574
0 4096.0 4.800000 9.600000
1 8192.0 8.727273 19.200000
2 16384.0 17.454545 38.400001
3 32768.0 38.400001 76.800002
4 65536.0 69.818181 127.999995
5 131072.0 139.636363 219.428568
6 262144.0 219.428568 341.333321
7 524288.0 341.333321 472.615390
8 1048576.0 472.615390 614.400016
9 2097152.0 614.400016 702.171410
10 4194304.0 712.347810 780.190482
11 8388608.0 774.047204 812.429770
12 16777216.0 809.086412 833.084721
13 33554432.0 829.569620 842.004273
14 67108864.0 840.205105 848.362445
15 134217728.0 845.625825 850.656574
@@ -255,7 +260,7 @@ We can now run the decorated function above. Pass `print_data=True` to see the p
.. rst-class:: sphx-glr-timing
**Total running time of the script:** ( 1 minutes 42.600 seconds)
**Total running time of the script:** ( 1 minutes 42.917 seconds)
.. _sphx_glr_download_getting-started_tutorials_01-vector-add.py:

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

@@ -278,17 +278,17 @@ We will then compare its performance against (1) :code:`torch.softmax` and (2) t
softmax-performance:
N Triton Torch (native) Torch (jit)
0 256.0 512.000001 546.133347 186.181817
1 384.0 614.400016 585.142862 153.600004
2 512.0 655.360017 606.814814 154.566038
3 640.0 706.206879 640.000002 158.759699
4 768.0 722.823517 664.216187 162.754967
0 256.0 512.000001 546.133347 190.511628
1 384.0 585.142862 585.142862 151.703707
2 512.0 655.360017 585.142849 154.566038
3 640.0 682.666684 640.000002 160.000000
4 768.0 722.823517 646.736871 163.839992
.. ... ... ... ...
93 12160.0 814.058574 406.179533 198.834951
94 12288.0 814.111783 415.661740 199.096718
95 12416.0 814.163950 412.149375 198.655991
96 12544.0 814.214963 412.971190 198.913776
97 12672.0 814.265046 411.679167 198.971549
93 12160.0 814.058574 405.755985 198.834951
94 12288.0 814.111783 415.222812 199.197579
95 12416.0 814.163950 412.149375 198.854847
96 12544.0 814.214963 412.971190 199.012395
97 12672.0 814.265046 412.097543 199.167004
[98 rows x 4 columns]
@@ -306,7 +306,7 @@ In the above plot, we can see that:
.. rst-class:: sphx-glr-timing
**Total running time of the script:** ( 3 minutes 19.345 seconds)
**Total running time of the script:** ( 3 minutes 27.571 seconds)
.. _sphx_glr_download_getting-started_tutorials_02-fused-softmax.py:

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

@@ -459,36 +459,36 @@ We can now compare the performance of our kernel against that of cuBLAS. Here we
matmul-performance:
M cuBLAS ... Triton Triton (+ LeakyReLU)
0 256.0 2.730667 ... 3.276800 2.978909
1 384.0 7.372800 ... 8.507077 7.899428
2 512.0 14.563555 ... 16.384000 16.384000
1 384.0 7.372800 ... 7.899428 8.507077
2 512.0 14.563555 ... 15.420235 15.420235
3 640.0 22.260869 ... 24.380953 24.380953
4 768.0 32.768000 ... 34.028308 34.028308
5 896.0 39.025776 ... 40.140799 39.025776
6 1024.0 49.932191 ... 52.428801 52.428801
7 1152.0 45.242181 ... 46.656000 46.656000
8 1280.0 51.200001 ... 56.888887 56.888887
9 1408.0 64.138541 ... 67.305878 66.485074
10 1536.0 80.430545 ... 79.526831 78.643199
11 1664.0 62.929456 ... 62.492442 61.636381
12 1792.0 72.512412 ... 72.512412 72.047592
13 1920.0 69.467336 ... 70.172588 69.818184
14 2048.0 73.262953 ... 76.608294 76.608294
15 2176.0 83.500614 ... 85.998493 85.632545
16 2304.0 68.643310 ... 77.057651 76.319081
17 2432.0 71.305746 ... 85.393507 85.134737
18 2560.0 78.019048 ... 80.709358 81.108913
19 2688.0 83.004501 ... 89.464755 89.254248
20 2816.0 80.767055 ... 83.552120 82.602666
21 2944.0 81.298583 ... 82.237674 81.967162
22 3072.0 81.707223 ... 88.473602 87.516392
23 3200.0 84.377059 ... 94.955488 94.674553
24 3328.0 83.034941 ... 84.695641 83.905938
25 3456.0 81.108217 ... 85.133652 81.029251
26 3584.0 86.623693 ... 98.483450 98.160909
27 3712.0 81.682211 ... 88.404730 84.017953
28 3840.0 82.592983 ... 92.006659 85.169042
29 3968.0 90.791620 ... 86.236000 90.522206
30 4096.0 86.369197 ... 85.001726 91.304576
4 768.0 32.768000 ... 35.389441 34.028308
5 896.0 37.971025 ... 40.140799 39.025776
6 1024.0 49.932191 ... 53.773130 53.773130
7 1152.0 45.242181 ... 48.161033 47.396572
8 1280.0 51.200001 ... 57.690139 57.690139
9 1408.0 64.138541 ... 69.009825 67.305878
10 1536.0 79.526831 ... 80.430545 79.526831
11 1664.0 63.372618 ... 63.372618 62.929456
12 1792.0 72.983276 ... 63.142831 63.142831
13 1920.0 69.120002 ... 71.626943 70.892307
14 2048.0 73.584279 ... 78.033565 77.672296
15 2176.0 83.500614 ... 87.115360 86.739860
16 2304.0 68.446623 ... 77.810656 77.307030
17 2432.0 71.125224 ... 75.522751 75.320281
18 2560.0 77.833728 ... 81.715711 81.512437
19 2688.0 83.737433 ... 90.966561 90.532356
20 2816.0 79.443003 ... 82.916747 84.035084
21 2944.0 81.832567 ... 83.198715 82.102191
22 3072.0 81.121923 ... 88.197981 87.381335
23 3200.0 83.116885 ... 96.385543 96.096095
24 3328.0 82.939284 ... 84.895397 85.096096
25 3456.0 77.745004 ... 86.596744 84.332184
26 3584.0 86.540320 ... 91.380335 97.416461
27 3712.0 85.748791 ... 88.404730 86.341700
28 3840.0 84.874902 ... 93.247896 85.300426
29 3968.0 92.372393 ... 80.015697 78.220472
30 4096.0 93.990003 ... 93.206754 86.036145
[31 rows x 5 columns]
@@ -498,7 +498,7 @@ We can now compare the performance of our kernel against that of cuBLAS. Here we
.. rst-class:: sphx-glr-timing
**Total running time of the script:** ( 6 minutes 0.645 seconds)
**Total running time of the script:** ( 6 minutes 48.521 seconds)
.. _sphx_glr_download_getting-started_tutorials_03-matrix-multiplication.py:

2
master/_sources/getting-started/tutorials/04-low-memory-dropout.rst.txt
View File

@@ -240,7 +240,7 @@ References
.. rst-class:: sphx-glr-timing
**Total running time of the script:** ( 0 minutes 0.482 seconds)
**Total running time of the script:** ( 0 minutes 0.327 seconds)
.. _sphx_glr_download_getting-started_tutorials_04-low-memory-dropout.py:

62
master/_sources/getting-started/tutorials/05-layer-norm.rst.txt
View File

@@ -38,36 +38,36 @@ Layer Normalization
layer-norm-backward:
N Triton Torch Apex
0 1024.0 311.088617 98.303995 303.407414
1 1536.0 347.773587 134.050910 341.333333
2 2048.0 420.102553 161.684218 325.509933
3 2560.0 455.111129 181.238943 326.808501
4 3072.0 511.999982 191.999993 317.793096
5 3584.0 547.872604 207.768111 310.527060
6 4096.0 564.965515 219.919464 295.207204
7 4608.0 504.986315 232.825259 291.799469
8 5120.0 527.381977 242.845844 288.450695
9 5632.0 542.843364 243.107920 288.820505
10 6144.0 546.133354 248.661056 286.879370
11 6656.0 532.479975 256.000009 285.767438
12 7168.0 505.976473 260.654538 286.242939
13 7680.0 481.253256 262.190612 278.429013
14 8192.0 463.698115 267.130429 284.939124
15 8704.0 417.791980 267.815384 284.987724
16 9216.0 432.845409 272.394084 288.751954
17 9728.0 439.683593 280.278512 289.667485
18 10240.0 448.467168 286.433562 290.153487
19 10752.0 425.821771 247.172406 290.922209
20 11264.0 427.071098 245.760001 286.676558
21 11776.0 423.724129 249.667843 288.981596
22 12288.0 419.504980 254.673582 294.323369
23 12800.0 413.458944 253.674644 288.180121
24 13312.0 412.242569 252.859526 289.916513
25 13824.0 405.594132 257.190689 292.056329
26 14336.0 394.568805 254.297107 286.719986
27 14848.0 386.498925 257.665934 289.246765
28 15360.0 373.117425 257.970599 286.211174
29 15872.0 371.274849 261.806182 289.899545
0 1024.0 356.173905 99.497980 315.076934
1 1536.0 405.098894 134.050910 344.523365
2 2048.0 486.653476 159.067963 321.254900
3 2560.0 458.507457 182.314537 326.808501
4 3072.0 515.580429 191.501303 319.168834
5 3584.0 551.384634 207.768111 307.199992
6 4096.0 568.231237 220.907859 293.444785
7 4608.0 502.690905 232.336141 290.267724
8 5120.0 527.381977 243.326731 287.102804
9 5632.0 540.671974 244.426754 291.310338
10 6144.0 548.163546 251.202731 288.000001
11 6656.0 532.479975 255.590406 286.279570
12 7168.0 510.480705 253.734520 277.919225
13 7680.0 487.619051 266.743841 284.884090
14 8192.0 468.114289 258.694729 278.481578
15 8704.0 415.300208 267.472468 284.987724
16 9216.0 429.483477 272.394084 290.077383
17 9728.0 438.033784 280.278512 288.950501
18 10240.0 443.610086 287.102804 290.153487
19 10752.0 426.525614 246.699797 290.267711
20 11264.0 427.071098 245.536784 286.069848
21 11776.0 418.702211 249.447482 288.981596
22 12288.0 414.784810 254.673582 294.323369
23 12800.0 410.146863 254.094291 288.180121
24 13312.0 409.599999 252.161013 289.129403
25 13824.0 404.112047 257.190689 291.799461
26 14336.0 395.930964 256.000002 289.129416
27 14848.0 385.662341 257.479779 288.777966
28 15360.0 380.433442 258.332158 286.656296
29 15872.0 372.363640 261.806182 290.562936
@@ -339,7 +339,7 @@ Layer Normalization
.. rst-class:: sphx-glr-timing
**Total running time of the script:** ( 2 minutes 12.550 seconds)
**Total running time of the script:** ( 2 minutes 14.583 seconds)
.. _sphx_glr_download_getting-started_tutorials_05-layer-norm.py:

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

@@ -5,16 +5,16 @@
Computation times
=================
**13:15.622** total execution time for **getting-started_tutorials** files:
**14:13.919** total execution time for **getting-started_tutorials** files:
+---------------------------------------------------------------------------------------------------------+-----------+--------+
| :ref:`sphx_glr_getting-started_tutorials_03-matrix-multiplication.py` (``03-matrix-multiplication.py``) | 06:00.645 | 0.0 MB |
| :ref:`sphx_glr_getting-started_tutorials_03-matrix-multiplication.py` (``03-matrix-multiplication.py``) | 06:48.521 | 0.0 MB |
+---------------------------------------------------------------------------------------------------------+-----------+--------+
| :ref:`sphx_glr_getting-started_tutorials_02-fused-softmax.py` (``02-fused-softmax.py``) | 03:19.345 | 0.0 MB |
| :ref:`sphx_glr_getting-started_tutorials_02-fused-softmax.py` (``02-fused-softmax.py``) | 03:27.571 | 0.0 MB |
+---------------------------------------------------------------------------------------------------------+-----------+--------+
| :ref:`sphx_glr_getting-started_tutorials_05-layer-norm.py` (``05-layer-norm.py``) | 02:12.550 | 0.0 MB |
| :ref:`sphx_glr_getting-started_tutorials_05-layer-norm.py` (``05-layer-norm.py``) | 02:14.583 | 0.0 MB |
+---------------------------------------------------------------------------------------------------------+-----------+--------+
| :ref:`sphx_glr_getting-started_tutorials_01-vector-add.py` (``01-vector-add.py``) | 01:42.600 | 0.0 MB |
| :ref:`sphx_glr_getting-started_tutorials_01-vector-add.py` (``01-vector-add.py``) | 01:42.917 | 0.0 MB |
+---------------------------------------------------------------------------------------------------------+-----------+--------+
| :ref:`sphx_glr_getting-started_tutorials_04-low-memory-dropout.py` (``04-low-memory-dropout.py``) | 00:00.482 | 0.0 MB |
| :ref:`sphx_glr_getting-started_tutorials_04-low-memory-dropout.py` (``04-low-memory-dropout.py``) | 00:00.327 | 0.0 MB |
+---------------------------------------------------------------------------------------------------------+-----------+--------+