diff --git a/master/.buildinfo b/master/.buildinfo index 27604d52c..1927b47f5 100644 --- a/master/.buildinfo +++ b/master/.buildinfo @@ -1,4 +1,4 @@ # Sphinx build info version 1 # This file hashes the configuration used when building these files. When it is not found, a full rebuild will be done. -config: 456fb2bf2c82c803cfdaa7bcb3e778a9 +config: aa051e1ca67e6f2658629dbe2c30cba0 tags: 645f666f9bcd5a90fca523b33c5a78b7 diff --git a/master/.doctrees/environment.pickle b/master/.doctrees/environment.pickle index 51eef543f..d57942e3b 100644 Binary files a/master/.doctrees/environment.pickle and b/master/.doctrees/environment.pickle differ diff --git a/master/.doctrees/getting-started/installation.doctree b/master/.doctrees/getting-started/installation.doctree index aeddca6d4..f5f6f0a71 100644 Binary files a/master/.doctrees/getting-started/installation.doctree and b/master/.doctrees/getting-started/installation.doctree differ diff --git a/master/.doctrees/getting-started/tutorials/01-vector-add.doctree b/master/.doctrees/getting-started/tutorials/01-vector-add.doctree index a5314eb19..b4fd24925 100644 Binary files a/master/.doctrees/getting-started/tutorials/01-vector-add.doctree and b/master/.doctrees/getting-started/tutorials/01-vector-add.doctree differ diff --git a/master/.doctrees/getting-started/tutorials/02-fused-softmax.doctree b/master/.doctrees/getting-started/tutorials/02-fused-softmax.doctree index 5caf7b6a1..3d287e9aa 100644 Binary files a/master/.doctrees/getting-started/tutorials/02-fused-softmax.doctree and b/master/.doctrees/getting-started/tutorials/02-fused-softmax.doctree differ diff --git a/master/.doctrees/getting-started/tutorials/03-matrix-multiplication.doctree b/master/.doctrees/getting-started/tutorials/03-matrix-multiplication.doctree index 9e2c12ece..9aed8a142 100644 Binary files a/master/.doctrees/getting-started/tutorials/03-matrix-multiplication.doctree and b/master/.doctrees/getting-started/tutorials/03-matrix-multiplication.doctree differ diff --git a/master/.doctrees/getting-started/tutorials/04-low-memory-dropout.doctree b/master/.doctrees/getting-started/tutorials/04-low-memory-dropout.doctree index 19a237748..70f20150d 100644 Binary files a/master/.doctrees/getting-started/tutorials/04-low-memory-dropout.doctree and b/master/.doctrees/getting-started/tutorials/04-low-memory-dropout.doctree differ diff --git a/master/.doctrees/getting-started/tutorials/05-layer-norm.doctree b/master/.doctrees/getting-started/tutorials/05-layer-norm.doctree index 266e9d4a3..62eb6bb44 100644 Binary files a/master/.doctrees/getting-started/tutorials/05-layer-norm.doctree and b/master/.doctrees/getting-started/tutorials/05-layer-norm.doctree differ diff --git a/master/.doctrees/getting-started/tutorials/index.doctree b/master/.doctrees/getting-started/tutorials/index.doctree index 7b5790de5..23adb18bc 100644 Binary files a/master/.doctrees/getting-started/tutorials/index.doctree and b/master/.doctrees/getting-started/tutorials/index.doctree differ diff --git a/master/.doctrees/getting-started/tutorials/sg_execution_times.doctree b/master/.doctrees/getting-started/tutorials/sg_execution_times.doctree index 0c2a47d6e..b38b1e4c7 100644 Binary files a/master/.doctrees/getting-started/tutorials/sg_execution_times.doctree and b/master/.doctrees/getting-started/tutorials/sg_execution_times.doctree differ diff --git a/master/.doctrees/index.doctree b/master/.doctrees/index.doctree index 1c542e1eb..f5a49c5ee 100644 Binary files a/master/.doctrees/index.doctree and b/master/.doctrees/index.doctree differ diff --git a/master/.doctrees/programming-guide/chapter-1/introduction.doctree b/master/.doctrees/programming-guide/chapter-1/introduction.doctree index 88eca63f5..e157a4cfa 100644 Binary files a/master/.doctrees/programming-guide/chapter-1/introduction.doctree and b/master/.doctrees/programming-guide/chapter-1/introduction.doctree differ diff --git a/master/.doctrees/programming-guide/chapter-2/related-work.doctree b/master/.doctrees/programming-guide/chapter-2/related-work.doctree index b12db3964..6491fe12b 100644 Binary files a/master/.doctrees/programming-guide/chapter-2/related-work.doctree and b/master/.doctrees/programming-guide/chapter-2/related-work.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.Config.doctree b/master/.doctrees/python-api/generated/triton.Config.doctree index e4c4f9de6..d8f49e3ca 100644 Binary files a/master/.doctrees/python-api/generated/triton.Config.doctree and b/master/.doctrees/python-api/generated/triton.Config.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.autotune.doctree b/master/.doctrees/python-api/generated/triton.autotune.doctree index 3e769328b..dcd3f1c8f 100644 Binary files a/master/.doctrees/python-api/generated/triton.autotune.doctree and b/master/.doctrees/python-api/generated/triton.autotune.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.heuristics.doctree b/master/.doctrees/python-api/generated/triton.heuristics.doctree index 5bfe88276..13e83b997 100644 Binary files a/master/.doctrees/python-api/generated/triton.heuristics.doctree and b/master/.doctrees/python-api/generated/triton.heuristics.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.jit.doctree b/master/.doctrees/python-api/generated/triton.jit.doctree index 04fa26a7a..d1bc2dab8 100644 Binary files a/master/.doctrees/python-api/generated/triton.jit.doctree and b/master/.doctrees/python-api/generated/triton.jit.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.arange.doctree b/master/.doctrees/python-api/generated/triton.language.arange.doctree index fa7329484..4eb5b588f 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.arange.doctree and b/master/.doctrees/python-api/generated/triton.language.arange.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.atomic_add.doctree b/master/.doctrees/python-api/generated/triton.language.atomic_add.doctree index 469064874..aca062b83 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.atomic_add.doctree and b/master/.doctrees/python-api/generated/triton.language.atomic_add.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.atomic_cas.doctree b/master/.doctrees/python-api/generated/triton.language.atomic_cas.doctree index 2b7a00deb..2da19a0fc 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.atomic_cas.doctree and b/master/.doctrees/python-api/generated/triton.language.atomic_cas.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.atomic_max.doctree b/master/.doctrees/python-api/generated/triton.language.atomic_max.doctree index 468d0938d..7fbf81859 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.atomic_max.doctree and b/master/.doctrees/python-api/generated/triton.language.atomic_max.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.atomic_min.doctree b/master/.doctrees/python-api/generated/triton.language.atomic_min.doctree index 562b8e5fd..7ba78acbb 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.atomic_min.doctree and b/master/.doctrees/python-api/generated/triton.language.atomic_min.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.atomic_xchg.doctree b/master/.doctrees/python-api/generated/triton.language.atomic_xchg.doctree index 680ab77e4..751bfedec 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.atomic_xchg.doctree and b/master/.doctrees/python-api/generated/triton.language.atomic_xchg.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.broadcast_to.doctree b/master/.doctrees/python-api/generated/triton.language.broadcast_to.doctree index 210a06f9f..72a9221ca 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.broadcast_to.doctree and b/master/.doctrees/python-api/generated/triton.language.broadcast_to.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.cos.doctree b/master/.doctrees/python-api/generated/triton.language.cos.doctree index 7ef13f742..b5d01a9ab 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.cos.doctree and b/master/.doctrees/python-api/generated/triton.language.cos.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.dot.doctree b/master/.doctrees/python-api/generated/triton.language.dot.doctree index 8a65688dd..d8b587815 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.dot.doctree and b/master/.doctrees/python-api/generated/triton.language.dot.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.exp.doctree b/master/.doctrees/python-api/generated/triton.language.exp.doctree index 7009e36c7..4a3ccb9c8 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.exp.doctree and b/master/.doctrees/python-api/generated/triton.language.exp.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.load.doctree b/master/.doctrees/python-api/generated/triton.language.load.doctree index c2e86d128..ad04b6c61 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.load.doctree and b/master/.doctrees/python-api/generated/triton.language.load.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.log.doctree b/master/.doctrees/python-api/generated/triton.language.log.doctree index a71afb324..acc41de3f 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.log.doctree and b/master/.doctrees/python-api/generated/triton.language.log.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.max.doctree b/master/.doctrees/python-api/generated/triton.language.max.doctree index 9122d4278..5d70e4566 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.max.doctree and b/master/.doctrees/python-api/generated/triton.language.max.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.maximum.doctree b/master/.doctrees/python-api/generated/triton.language.maximum.doctree index 983b40b14..e4ccfc22d 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.maximum.doctree and b/master/.doctrees/python-api/generated/triton.language.maximum.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.min.doctree b/master/.doctrees/python-api/generated/triton.language.min.doctree index dec63115c..57d51cf94 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.min.doctree and b/master/.doctrees/python-api/generated/triton.language.min.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.minimum.doctree b/master/.doctrees/python-api/generated/triton.language.minimum.doctree index 8bb5a5fba..72602e89c 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.minimum.doctree and b/master/.doctrees/python-api/generated/triton.language.minimum.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.multiple_of.doctree b/master/.doctrees/python-api/generated/triton.language.multiple_of.doctree index ec05b5680..e58ab59fb 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.multiple_of.doctree and b/master/.doctrees/python-api/generated/triton.language.multiple_of.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.num_programs.doctree b/master/.doctrees/python-api/generated/triton.language.num_programs.doctree index cb56731b6..c944c51c3 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.num_programs.doctree and b/master/.doctrees/python-api/generated/triton.language.num_programs.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.program_id.doctree b/master/.doctrees/python-api/generated/triton.language.program_id.doctree index 3813a4aff..9ec57f95a 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.program_id.doctree and b/master/.doctrees/python-api/generated/triton.language.program_id.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.rand.doctree b/master/.doctrees/python-api/generated/triton.language.rand.doctree index 3cbd1e6cc..35d7cf58b 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.rand.doctree and b/master/.doctrees/python-api/generated/triton.language.rand.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.randint.doctree b/master/.doctrees/python-api/generated/triton.language.randint.doctree index b264b4eb7..3b66880ed 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.randint.doctree and b/master/.doctrees/python-api/generated/triton.language.randint.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.randint4x.doctree b/master/.doctrees/python-api/generated/triton.language.randint4x.doctree index 441ec42b1..3b35387a1 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.randint4x.doctree and b/master/.doctrees/python-api/generated/triton.language.randint4x.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.randn.doctree b/master/.doctrees/python-api/generated/triton.language.randn.doctree index a5f67d2a5..3ad74e17c 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.randn.doctree and b/master/.doctrees/python-api/generated/triton.language.randn.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.ravel.doctree b/master/.doctrees/python-api/generated/triton.language.ravel.doctree index 574624fd6..dea119335 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.ravel.doctree and b/master/.doctrees/python-api/generated/triton.language.ravel.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.reshape.doctree b/master/.doctrees/python-api/generated/triton.language.reshape.doctree index 591f75aa0..e6cd64dff 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.reshape.doctree and b/master/.doctrees/python-api/generated/triton.language.reshape.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.sigmoid.doctree b/master/.doctrees/python-api/generated/triton.language.sigmoid.doctree index 8d1272a3b..15bf1a627 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.sigmoid.doctree and b/master/.doctrees/python-api/generated/triton.language.sigmoid.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.sin.doctree b/master/.doctrees/python-api/generated/triton.language.sin.doctree index e4305d99c..b1faad4ab 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.sin.doctree and b/master/.doctrees/python-api/generated/triton.language.sin.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.softmax.doctree b/master/.doctrees/python-api/generated/triton.language.softmax.doctree index f275dd969..59a86d78e 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.softmax.doctree and b/master/.doctrees/python-api/generated/triton.language.softmax.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.sqrt.doctree b/master/.doctrees/python-api/generated/triton.language.sqrt.doctree index 715a7c288..0cd05fd0c 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.sqrt.doctree and b/master/.doctrees/python-api/generated/triton.language.sqrt.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.store.doctree b/master/.doctrees/python-api/generated/triton.language.store.doctree index d99c5ca9e..865c218ba 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.store.doctree and b/master/.doctrees/python-api/generated/triton.language.store.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.sum.doctree b/master/.doctrees/python-api/generated/triton.language.sum.doctree index 0c7a140e1..8dd1e85ec 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.sum.doctree and b/master/.doctrees/python-api/generated/triton.language.sum.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.where.doctree b/master/.doctrees/python-api/generated/triton.language.where.doctree index b3b81797f..d6fef1e44 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.where.doctree and b/master/.doctrees/python-api/generated/triton.language.where.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.language.zeros.doctree b/master/.doctrees/python-api/generated/triton.language.zeros.doctree index 8714bbcf4..db39c8bdf 100644 Binary files a/master/.doctrees/python-api/generated/triton.language.zeros.doctree and b/master/.doctrees/python-api/generated/triton.language.zeros.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.testing.Benchmark.doctree b/master/.doctrees/python-api/generated/triton.testing.Benchmark.doctree index c2abe1815..9e83ee96f 100644 Binary files a/master/.doctrees/python-api/generated/triton.testing.Benchmark.doctree and b/master/.doctrees/python-api/generated/triton.testing.Benchmark.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.testing.do_bench.doctree b/master/.doctrees/python-api/generated/triton.testing.do_bench.doctree index ddeed4fd9..1a85a490a 100644 Binary files a/master/.doctrees/python-api/generated/triton.testing.do_bench.doctree and b/master/.doctrees/python-api/generated/triton.testing.do_bench.doctree differ diff --git a/master/.doctrees/python-api/generated/triton.testing.perf_report.doctree b/master/.doctrees/python-api/generated/triton.testing.perf_report.doctree index a7b1f6705..b0905f553 100644 Binary files a/master/.doctrees/python-api/generated/triton.testing.perf_report.doctree and b/master/.doctrees/python-api/generated/triton.testing.perf_report.doctree differ diff --git a/master/.doctrees/python-api/triton.doctree b/master/.doctrees/python-api/triton.doctree index 3661712c4..7c643b2a7 100644 Binary files a/master/.doctrees/python-api/triton.doctree and b/master/.doctrees/python-api/triton.doctree differ diff --git a/master/.doctrees/python-api/triton.language.doctree b/master/.doctrees/python-api/triton.language.doctree index aa9699b65..5039c0111 100644 Binary files a/master/.doctrees/python-api/triton.language.doctree and b/master/.doctrees/python-api/triton.language.doctree differ diff --git a/master/.doctrees/python-api/triton.testing.doctree b/master/.doctrees/python-api/triton.testing.doctree index 0ce56bf0b..86c80bf9f 100644 Binary files a/master/.doctrees/python-api/triton.testing.doctree and b/master/.doctrees/python-api/triton.testing.doctree differ diff --git a/master/_downloads/62d97d49a32414049819dd8bb8378080/01-vector-add.py b/master/_downloads/62d97d49a32414049819dd8bb8378080/01-vector-add.py index d684106f1..51de7ac6c 100644 --- a/master/_downloads/62d97d49a32414049819dd8bb8378080/01-vector-add.py +++ b/master/_downloads/62d97d49a32414049819dd8bb8378080/01-vector-add.py @@ -24,9 +24,11 @@ def add_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 @@ -45,6 +47,7 @@ def add_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()) # %% @@ -53,6 +56,8 @@ def add_kernel( 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 @@ -65,7 +70,7 @@ def add(x: torch.Tensor, y: torch.Tensor): # - 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 diff --git a/master/_downloads/662999063954282841dc90b8945f85ce/tutorials_jupyter.zip b/master/_downloads/662999063954282841dc90b8945f85ce/tutorials_jupyter.zip index b1bd1383f..3ececce53 100644 Binary files a/master/_downloads/662999063954282841dc90b8945f85ce/tutorials_jupyter.zip and b/master/_downloads/662999063954282841dc90b8945f85ce/tutorials_jupyter.zip differ diff --git a/master/_downloads/763344228ae6bc253ed1a6cf586aa30d/tutorials_python.zip b/master/_downloads/763344228ae6bc253ed1a6cf586aa30d/tutorials_python.zip index 5dc522139..95b807c77 100644 Binary files a/master/_downloads/763344228ae6bc253ed1a6cf586aa30d/tutorials_python.zip and b/master/_downloads/763344228ae6bc253ed1a6cf586aa30d/tutorials_python.zip differ diff --git a/master/_downloads/f191ee1e78dc52eb5f7cba88f71cef2f/01-vector-add.ipynb b/master/_downloads/f191ee1e78dc52eb5f7cba88f71cef2f/01-vector-add.ipynb index a88ec6569..45bef9d0b 100644 --- a/master/_downloads/f191ee1e78dc52eb5f7cba88f71cef2f/01-vector-add.ipynb +++ b/master/_downloads/f191ee1e78dc52eb5f7cba88f71cef2f/01-vector-add.ipynb @@ -33,7 +33,7 @@ }, "outputs": [], "source": [ - "import torch\n\nimport triton\nimport triton.language as tl\n\n\n@triton.jit\ndef add_kernel(\n x_ptr, # *Pointer* to first input vector\n y_ptr, # *Pointer* to second input vector\n output_ptr, # *Pointer* to output vector\n n_elements, # Size of the vector\n BLOCK_SIZE: tl.constexpr, # Number of elements each program should process\n # NOTE: `constexpr` so it can be used as a shape value\n):\n # There are multiple 'program's processing different data. We identify which program\n # we are here\n pid = tl.program_id(axis=0) # We use a 1D launch grid so axis is 0\n # This program will process inputs that are offset from the initial data.\n # for instance, if you had a vector of length 256 and block_size of 64, the programs\n # would each access the elements [0:64, 64:128, 128:192, 192:256].\n # Note that offsets is a list of pointers\n block_start = pid * BLOCK_SIZE\n offsets = block_start + tl.arange(0, BLOCK_SIZE)\n # Create a mask to guard memory operations against out-of-bounds accesses\n mask = offsets < n_elements\n # Load x and y from DRAM, masking out any extra elements in case the input is not a\n # multiple of the block size\n x = tl.load(x_ptr + offsets, mask=mask)\n y = tl.load(y_ptr + offsets, mask=mask)\n output = x + y\n # Write x + y back to DRAM\n tl.store(output_ptr + offsets, output, mask=mask)" + "import torch\n\nimport triton\nimport triton.language as tl\n\n\n@triton.jit\ndef add_kernel(\n x_ptr, # *Pointer* to first input vector\n y_ptr, # *Pointer* to second input vector\n output_ptr, # *Pointer* to output vector\n n_elements, # Size of the vector\n time_start_ptr, time_end_ptr,\n BLOCK_SIZE: tl.constexpr, # Number of elements each program should process\n # NOTE: `constexpr` so it can be used as a shape value\n):\n tl.atomic_min(time_start_ptr, tl.clock())\n # There are multiple 'program's processing different data. We identify which program\n # we are here\n pid = tl.program_id(axis=0) # We use a 1D launch grid so axis is 0\n # This program will process inputs that are offset from the initial data.\n # for instance, if you had a vector of length 256 and block_size of 64, the programs\n # would each access the elements [0:64, 64:128, 128:192, 192:256].\n # Note that offsets is a list of pointers\n block_start = pid * BLOCK_SIZE\n offsets = block_start + tl.arange(0, BLOCK_SIZE)\n # Create a mask to guard memory operations against out-of-bounds accesses\n mask = offsets < n_elements\n # Load x and y from DRAM, masking out any extra elements in case the input is not a\n # multiple of the block size\n x = tl.load(x_ptr + offsets, mask=mask)\n y = tl.load(y_ptr + offsets, mask=mask)\n output = x + y\n # Write x + y back to DRAM\n tl.store(output_ptr + offsets, output, mask=mask)\n tl.atomic_max(time_end_ptr, tl.clock())" ] }, { @@ -51,7 +51,7 @@ }, "outputs": [], "source": [ - "def add(x: torch.Tensor, y: torch.Tensor):\n # We need to preallocate the output\n output = torch.empty_like(x)\n assert x.is_cuda and y.is_cuda and output.is_cuda\n n_elements = output.numel()\n # The SPMD launch grid denotes the number of kernel instances that run in parallel.\n # It is analogous to CUDA launch grids. It can be either Tuple[int], or Callable(metaparameters) -> Tuple[int]\n # In this case, we use a 1D grid where the size is the number of blocks\n grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']),)\n # NOTE:\n # - each torch.tensor object is implicitly converted into a pointer to its first element.\n # - `triton.jit`'ed functions can be index with a launch grid to obtain a callable GPU kernel\n # - don't forget to pass meta-parameters as keywords arguments\n add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)\n # We return a handle to z but, since `torch.cuda.synchronize()` hasn't been called, the kernel is still\n # running asynchronously at this point.\n return output" + "def add(x: torch.Tensor, y: torch.Tensor):\n time_start = torch.zeros(1, dtype=torch.int64, device='cuda')\n time_end = torch.zeros(1, dtype=torch.int64, device='cuda')\n # We need to preallocate the output\n output = torch.empty_like(x)\n assert x.is_cuda and y.is_cuda and output.is_cuda\n n_elements = output.numel()\n # The SPMD launch grid denotes the number of kernel instances that run in parallel.\n # It is analogous to CUDA launch grids. It can be either Tuple[int], or Callable(metaparameters) -> Tuple[int]\n # In this case, we use a 1D grid where the size is the number of blocks\n grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']),)\n # NOTE:\n # - each torch.tensor object is implicitly converted into a pointer to its first element.\n # - `triton.jit`'ed functions can be index with a launch grid to obtain a callable GPU kernel\n # - don't forget to pass meta-parameters as keywords arguments\n add_kernel[grid](x, y, output, n_elements, time_start, time_end, BLOCK_SIZE=1024)\n # We return a handle to z but, since `torch.cuda.synchronize()` hasn't been called, the kernel is still\n # running asynchronously at this point.\n return output" ] }, { diff --git a/master/_images/sphx_glr_01-vector-add_001.png b/master/_images/sphx_glr_01-vector-add_001.png index 8ac13e126..8ff42448f 100644 Binary files a/master/_images/sphx_glr_01-vector-add_001.png and b/master/_images/sphx_glr_01-vector-add_001.png differ diff --git a/master/_images/sphx_glr_01-vector-add_thumb.png b/master/_images/sphx_glr_01-vector-add_thumb.png index bc9e49567..66e610bbe 100644 Binary files a/master/_images/sphx_glr_01-vector-add_thumb.png and b/master/_images/sphx_glr_01-vector-add_thumb.png differ diff --git a/master/_images/sphx_glr_02-fused-softmax_001.png b/master/_images/sphx_glr_02-fused-softmax_001.png index dbb633acf..4072661db 100644 Binary files a/master/_images/sphx_glr_02-fused-softmax_001.png and b/master/_images/sphx_glr_02-fused-softmax_001.png differ diff --git a/master/_images/sphx_glr_02-fused-softmax_thumb.png b/master/_images/sphx_glr_02-fused-softmax_thumb.png index c9f7fb12e..950f2da8d 100644 Binary files a/master/_images/sphx_glr_02-fused-softmax_thumb.png and b/master/_images/sphx_glr_02-fused-softmax_thumb.png differ diff --git a/master/_images/sphx_glr_03-matrix-multiplication_001.png b/master/_images/sphx_glr_03-matrix-multiplication_001.png index f4b9d30d7..a6f841bcc 100644 Binary files a/master/_images/sphx_glr_03-matrix-multiplication_001.png and b/master/_images/sphx_glr_03-matrix-multiplication_001.png differ diff --git a/master/_images/sphx_glr_03-matrix-multiplication_thumb.png b/master/_images/sphx_glr_03-matrix-multiplication_thumb.png index d0c48d331..255098687 100644 Binary files a/master/_images/sphx_glr_03-matrix-multiplication_thumb.png and b/master/_images/sphx_glr_03-matrix-multiplication_thumb.png differ diff --git a/master/_images/sphx_glr_05-layer-norm_001.png b/master/_images/sphx_glr_05-layer-norm_001.png index aa114592f..96ddbc96b 100644 Binary files a/master/_images/sphx_glr_05-layer-norm_001.png and b/master/_images/sphx_glr_05-layer-norm_001.png differ diff --git a/master/_images/sphx_glr_05-layer-norm_thumb.png b/master/_images/sphx_glr_05-layer-norm_thumb.png index 6012163b5..f0226894d 100644 Binary files a/master/_images/sphx_glr_05-layer-norm_thumb.png and b/master/_images/sphx_glr_05-layer-norm_thumb.png differ diff --git a/master/_sources/getting-started/tutorials/01-vector-add.rst.txt b/master/_sources/getting-started/tutorials/01-vector-add.rst.txt index 0a88e309b..2f9700c96 100644 --- a/master/_sources/getting-started/tutorials/01-vector-add.rst.txt +++ b/master/_sources/getting-started/tutorials/01-vector-add.rst.txt @@ -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: diff --git a/master/_sources/getting-started/tutorials/02-fused-softmax.rst.txt b/master/_sources/getting-started/tutorials/02-fused-softmax.rst.txt index 9326cf35a..d9160d19b 100644 --- a/master/_sources/getting-started/tutorials/02-fused-softmax.rst.txt +++ b/master/_sources/getting-started/tutorials/02-fused-softmax.rst.txt @@ -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: diff --git a/master/_sources/getting-started/tutorials/03-matrix-multiplication.rst.txt b/master/_sources/getting-started/tutorials/03-matrix-multiplication.rst.txt index 020e790a7..bf8b14884 100644 --- a/master/_sources/getting-started/tutorials/03-matrix-multiplication.rst.txt +++ b/master/_sources/getting-started/tutorials/03-matrix-multiplication.rst.txt @@ -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: diff --git a/master/_sources/getting-started/tutorials/04-low-memory-dropout.rst.txt b/master/_sources/getting-started/tutorials/04-low-memory-dropout.rst.txt index da96d2c09..4dd4cf9de 100644 --- a/master/_sources/getting-started/tutorials/04-low-memory-dropout.rst.txt +++ b/master/_sources/getting-started/tutorials/04-low-memory-dropout.rst.txt @@ -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: diff --git a/master/_sources/getting-started/tutorials/05-layer-norm.rst.txt b/master/_sources/getting-started/tutorials/05-layer-norm.rst.txt index 58bd93a57..b230a1826 100644 --- a/master/_sources/getting-started/tutorials/05-layer-norm.rst.txt +++ b/master/_sources/getting-started/tutorials/05-layer-norm.rst.txt @@ -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: diff --git a/master/_sources/getting-started/tutorials/sg_execution_times.rst.txt b/master/_sources/getting-started/tutorials/sg_execution_times.rst.txt index f0856eb7b..8975bcca2 100644 --- a/master/_sources/getting-started/tutorials/sg_execution_times.rst.txt +++ b/master/_sources/getting-started/tutorials/sg_execution_times.rst.txt @@ -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 | +---------------------------------------------------------------------------------------------------------+-----------+--------+ diff --git a/master/getting-started/tutorials/01-vector-add.html b/master/getting-started/tutorials/01-vector-add.html index f17f0ceb1..a778e7110 100644 --- a/master/getting-started/tutorials/01-vector-add.html +++ b/master/getting-started/tutorials/01-vector-add.html @@ -214,9 +214,11 @@ to download the full example code

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 @@ -235,11 +237,14 @@ to download the full example code

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())

Let’s also declare a helper function to (1) allocate the z tensor and (2) enqueue the above kernel with appropriate grid/block sizes.

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
@@ -252,7 +257,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
@@ -322,25 +327,25 @@ for different problem sizes.

 
Out:

 
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
 

 

-
Total running time of the script: ( 1 minutes  42.600 seconds)

+
Total running time of the script: ( 1 minutes  42.917 seconds)
-

Total running time of the script: ( 2 minutes 12.550 seconds)

+

Total running time of the script: ( 2 minutes 14.583 seconds)