Commit Graph

280 Commits

Author SHA1 Message Date
Philippe Tillet
3ab121dbdb [PYTHON] Added support for tuples (#116) 2021-07-27 12:38:49 -07:00
Philippe Tillet
b5dcac484d [CODEGEN] Small bugfix in atomic-add (#114) 2021-07-27 12:38:49 -07:00
Philippe Tillet
f81012a8cf [CODEGEN] Fixed atomic_add issue (#112)
* [CODEGEN] Fixed atomic_add issue

* [CODEGEN] Fixed liveness analysis bug for instructions that are not
DCE'd but have no users (e.g., atomic_cas)
2021-07-27 12:38:49 -07:00
Philippe Tillet
325ee38581 [PYTHON] Fixed bug in scoping mechanism (#111)
Inline functions didn't restore scope of parents. Also some control flow
structure still had the scoping semantics of C++
2021-07-27 12:38:49 -07:00
Philippe Tillet
9f30af76fb [GENERAL] Minor improvements: (#110)
* Load libcuda.so.1 if libcuda.so is not there. Error if both aren't
there.
* Support for multiple grad_to_none in triton.testing.do_bench
* Benchmark dataframe printed along with name
2021-07-27 12:38:49 -07:00
Philippe Tillet
288b4f7f58 [PYTHON] Added frontend to print sass using turingas disasm.py (#109) 2021-07-27 12:38:49 -07:00
Philippe Tillet
c91dd56a92 [CI] Made setup.py more backwards-compatible (#108) 2021-07-27 12:38:49 -07:00
daadaada
840d65d8c6 [CODEGEN] Clean up visit_mma884 (#107) 2021-07-27 12:38:49 -07:00
daadaada
967e629c0c [CODEGEN] Add a pass to prefetch operands of dot if applicable. (#105)
* update membar pass when data is double buffered

* Add instruction prefetch_s

* prefetch tests pass (except the 1 warp case)

* Fix the 1-warp bug

* Add back prefetch files

* Disable prefetch on a100

* Always add war barrier on sm>=80
2021-07-27 12:38:49 -07:00
Philippe Tillet
147675923e [triton-ops] Minor build improvements (#106) 2021-07-27 12:38:49 -07:00
Philippe Tillet
d10265f054 [CODEGEN] Bugfix for immediate offsets in inline PTX (#104) 2021-07-27 12:38:49 -07:00
Philippe Tillet
1e844ba78d [CODEGEN] Switching to predicated inline PTX for LDGs (#103) 2021-07-27 12:38:49 -07:00
Szymon Sidor
ac57812bdc [PYTHON] Update TensorWrapper with device attribute (#102)
bugfix `triton.reinterpret`
2021-07-27 12:38:49 -07:00
Philippe Tillet
840140bf26 [CODEGEN] Removed dedicated reassociate pass to merge it into LLVM isel (#101)
This massively simplifies implementation of `reassociate` and also fixes
a bunch of bug. The pass could still be improved, but can already be used
to generate constant pointer offsets in eg the matmul epilogue
2021-07-27 12:38:49 -07:00
Philippe Tillet
e16bee1a27 [CI] Now tagging the nightly build with the date of the latest commit (#100)
(instead of current date)
2021-07-27 12:38:49 -07:00
Philippe Tillet
6a9810ccf2 [codegen] small bugfix: (#97)
* Added fp32 -> fp8 for ConstantFP = 0
 * Added some more robust semantic check for atomic_add
2021-07-27 12:38:49 -07:00
Philippe Tillet
7355efa745 [LANG] Preliminary FP8 support (#96) 2021-07-27 12:38:49 -07:00
Philippe Tillet
4290be1ae8 [PYTHON] Various minor codegen fixes (#95) 2021-07-27 12:38:49 -07:00
Philippe Tillet
2b75158426 [PYTHON] Added atomic_add (#94) 2021-07-27 12:38:49 -07:00
Nora Belrose
d7f87929fa [PYTHON] Added shape & device checks for inputs to sparse matmul op (#93) 2021-07-27 12:38:49 -07:00
Philippe Tillet
bfc0a7587d [PYTHON] Renamed triton.core -> triton.language (#92) 2021-07-27 12:38:49 -07:00
Nora Belrose
41410012e8 [PYTHON] Fix crash in sparse matmul for tensors w/ less than 4 dims (#91) 2021-07-27 12:38:49 -07:00
daadaada
f6688372db [PYTHON] Allow triton.code_gen.Binary to print Triton-IR asm. (#89) 2021-07-27 12:38:49 -07:00
Nora Belrose
1112e2526e [PYTHON] Fix variable referenced before assignment error in blocksparse matmul backward (#90) 2021-07-27 12:38:49 -07:00
Philippe Tillet
d9112144b4 [PYTHON] Now triton.code_gen.Binary can print PTX and LLIR (#88) 2021-07-27 12:38:49 -07:00
Philippe Tillet
29e33e50b7 [DOCS] Updates and improvements (#87) 2021-07-27 12:38:49 -07:00
Philippe Tillet
39f4730305 Deprecation of Triton-C and Replacement by decorated Python functions (#86)
This PR implements a major overhaul of the frontend for Triton, and replaces Triton-C by a pure Python API in which kernels are defined as @triton.jit decorated functions. The documentation and tutorials have also been updated to accommodate these changes.

See documentations for more information on the new API
2021-07-27 12:38:49 -07:00
Philippe Tillet
1fdb465b71 [DOCS] Various improvements and typo fixes 2021-07-27 12:38:49 -07:00
Philippe Tillet
3f6ba1020d [CODEGEN] Make sure peephole is called before anything else in codegen 2021-07-27 12:38:49 -07:00
Philippe Tillet
167a2e4b1a [PYTHON] Fixed formatting issue in conv.c 2021-07-27 12:38:49 -07:00
Philippe Tillet
5ba5a77561 [BUILD] Remove compilation warnings 2021-07-27 12:38:49 -07:00
Philippe Tillet
b352bc79e3 [CI] Changed triton-nightly to --pre triton (#78)
The solution proposed in #77 can create namespace conflicts when triton and triton-nightly have both been pip installed. Therefore, this PR is moving nightly releases to pre-releases in the main triton index.
2021-07-27 12:38:49 -07:00
Philippe Tillet
2f80a98776 [BUILD] Added automatic nightly build releases to pip in CI; removed build-time dependence on LLVM and PyTorch (#77)
Recently there has been more and more report about installation issues:

    - Installing Triton before upgrading pytorch can create some issues because Triton uses some torch headers

    - llvm-10-dev not available on some platform; llvm-11-dev not available on e.g. Ubuntu.
    absence of nightly builds

This PR should fix all these issues. Some CMake tricks are used to download and install llvm at build time. Triton Python bindings were modified to remove dependence on pytorch ops. Midnight CI job added to generate binary wheels for all Triton version and update them on pypi's new triton-nightly project.

This PR will also make it very easy to use LLVM forks in the future for whatever needs we have.
2021-07-27 12:38:49 -07:00
Philippe Tillet
3ad0a4d7be [DOCS] Uncommented sphinx gallery 2021-07-27 12:38:49 -07:00
Philippe Tillet
a74919fa46 [DOCS] Improved index 2021-07-27 12:38:49 -07:00
Philippe Tillet
997e54e3bf [DOCS] Added non-tutorial documentation pages 2021-07-27 12:38:49 -07:00
Philippe Tillet
183878dce5 [DOCS] Added matrix multiplication tutorial 2021-07-27 12:38:49 -07:00
Philippe Tillet
f4fb209dad [DOCS] Removed pip installation instruction as version on Pip is not
up-to-date
2021-07-27 12:38:49 -07:00
Philippe Tillet
50e58d73db [DOCS] Improved plots in tutorials 2021-07-27 12:38:49 -07:00
Philippe Tillet
eacbb73968 [PYTHON] CUTLASS wrapper for fair benchmarks (#75)
Before this commit, the benchmarking infrastructure used heterogeneous protocols between library (e.g., CUTLASS uses a C++ binary that reports mean TFLOPS; torch and triton use python call and report 10th, 50th and 90th quantiles). For the sake of uniformity and fair benchmark practices, this PR adds a python wrapper for auto-tuned CUTLASS matrix multiplication. Benchmarks have been rewritten to use this wrapper with `triton.testing.do_bench` rather than system calls to CUTLASS profiler. Importantly, this also ensures that all the matmuls are done on the *same* input data which should stabilize clock across providers.
2021-07-27 12:38:49 -07:00
Philippe Tillet
d6f18742b1 [README] Now linking to the documentation 2021-07-27 12:38:49 -07:00
Philippe Tillet
58a5c87c53 [PYTHON] Made bench_blocksparse and bench_cross_entropy compatible
with the new performance report API
2021-07-27 12:38:49 -07:00
Philippe Tillet
5b9afaa688 [CODEGEN] Fixed bug that caused conditional operator to not always
properly mask load operations

Also includes minor improvement to benchmarking infrastructure
2021-07-27 12:38:49 -07:00
Philippe Tillet
d1d09566b1 [DOCS] Improved tutorials documentation 2021-07-27 12:38:49 -07:00
Philippe Tillet
85752037eb [PYTHON] Changed benchmarking strategy. Instead of enqueueing many
kernels before synchronizing, the kernels are now  enqueued one by one.

This makes it possible to clear the L2 cache before running the
workload, and also potentially collect some variance data for error bars
in plots
2021-07-27 12:38:49 -07:00
Philippe Tillet
92242ace2c [DOCS] Re-structured documentation hierarchy 2021-07-27 12:38:49 -07:00
Philippe Tillet
ca04da3575 [DOCS] Switched tutorials to Python and use Sphinx Gallery 2021-07-27 12:38:49 -07:00
Philippe Tillet
5172792543 [DOCS] Added .ipynb tutorials in docs 2021-07-27 12:38:49 -07:00
Philippe Tillet
0c13b8ff0e [DOCS] Updated and improved docs (#73) 2021-07-27 12:38:49 -07:00
Philippe Tillet
3ecf834a69 [PYTHON] Deleted 01-vector-add.py: it is an unnecessary duplicate of
01-vector-add.ipynb
2021-07-27 12:38:49 -07:00