These 31 commits are when the Protocol Buffers files have changed:
| Commit: | 8585ff7 | |
|---|---|---|
| Author: | Sergei Lebedev | |
| Committer: | jax authors | |
[triton] Added support for host-side TMA descriptors PiperOrigin-RevId: 963369899
| Commit: | 8e07bd4 | |
|---|---|---|
| Author: | Eusebio Durán Montaña | |
| Committer: | jax authors | |
Implement AOT compilation for autotuned JAX-Triton kernels. We now compile all the autotuning candidate configs during the instantiate phase (i.e. compile time in XLA), and later autotune them during instantiate. Some things worth noting: * We always run instantiate, and so we can always expect its state. During JIT compilations this just happens right before initialize. * Instantiate and Execute happen in the same process, so we don't need to serialize/deserialize protos, we can just use the pointers to the objects in the static cache. PiperOrigin-RevId: 963327724
| Commit: | 4ecb52d | |
|---|---|---|
| Author: | Sergei Lebedev | |
| Committer: | jax authors | |
Added support for host-side TMA descriptors Note that, unfortunately, `TensorDescriptor` in Gluon assumes PyTorch when validating the arguments. So, using it with jax-triton requires patching or removing `TensorDescriptor.__post_init__` as we do in tests. PiperOrigin-RevId: 962805739
| Commit: | 7df5ccc | |
|---|---|---|
| Author: | Sergei Lebedev | |
| Committer: | jax authors | |
[triton] Added support for device-side TMA descriptors Closes #38912 and #37482. PiperOrigin-RevId: 960937449
| Commit: | d35d53a | |
|---|---|---|
| Author: | Eusebio Durán Montaña | |
| Committer: | jax authors | |
Implement AOT compilation for non-autotuned JAX-Triton kernels. Previously, JAX-Triton kernels depended on a specific PTX version at runtime, which created backwards-compatibility constraints and increased startup latency when loading the binary. During FFI instantiation (compile time), compile non-autotuned kernels down to the CUBIN level so that executable binaries are ready immediately upon load. This removes the runtime PTX dependency, improves backwards compatibility, and reduces load times. PiperOrigin-RevId: 960276682
| Commit: | 5f3f0a1 | |
|---|---|---|
| Author: | Eusebio Durán Montaña | |
| Committer: | jax authors | |
Add AOT tests for jax-triton kernels in XLA Created a py script that generates the HLO of the kernel, and added a test on the XLA side that compiles, and runs an AOT binary using that kernel. I didn't want to leak out XLA test rules in jaxlib so I kept the jax parts in the jax directory, and the XLA test in the XLA dir. PiperOrigin-RevId: 958348877
| Commit: | 8bebd73 | |
|---|---|---|
| Author: | Eusebio Durán Montaña | |
| Committer: | jax authors | |
Implement AOT compilation for non-autotuned JAX-Triton kernels. Previously, JAX-Triton kernels depended on a specific PTX version at runtime, which created backwards-compatibility constraints and increased startup latency when loading the binary. During FFI instantiation (compile time), compile non-autotuned kernels down to the CUBIN level so that executable binaries are ready immediately upon load. This removes the runtime PTX dependency, improves backwards compatibility, and reduces load times. PiperOrigin-RevId: 958260793
| Commit: | 62f82fb | |
|---|---|---|
| Author: | Eusebio Durán Montaña | |
| Committer: | jax authors | |
Implement AOT compilation for autotuned JAX-Triton kernels. We now compile all the autotuning candidate configs during the instantiate phase (i.e. compile time in XLA), and later autotune them during instantiate. PiperOrigin-RevId: 959692915
| Commit: | 53d03c0 | |
|---|---|---|
| Author: | Eusebio Durán Montaña | |
| Committer: | jax authors | |
Add AOT tests for jax-triton kernels in XLA Created a py script that generates the HLO of the kernel, and added a test on the XLA side that compiles, and runs an AOT binary using that kernel. I didn't want to leak out XLA test rules in jaxlib so I kept the jax parts in the jax directory, and the XLA test in the XLA dir. PiperOrigin-RevId: 958348877
| Commit: | 19080ed | |
|---|---|---|
| Author: | Eusebio Durán Montaña | |
| Committer: | jax authors | |
Implement AOT compilation for non-autotuned JAX-Triton kernels. Previously, JAX-Triton kernels depended on a specific PTX version at runtime, which created backwards-compatibility constraints and increased startup latency when loading the binary. During FFI instantiation (compile time), compile non-autotuned kernels down to the CUBIN level so that executable binaries are ready immediately upon load. This removes the runtime PTX dependency, improves backwards compatibility, and reduces load times. PiperOrigin-RevId: 958260793
| Commit: | db9080b | |
|---|---|---|
| Author: | Eusebio Durán Montaña | |
| Committer: | jax authors | |
Implement AOT compilation for non-autotuned JAX-Triton kernels. Previously, JAX-Triton kernels depended on a specific PTX version at runtime, which created backwards-compatibility constraints and increased startup latency when loading the binary. During FFI instantiation (compile time), compile non-autotuned kernels down to the CUBIN level so that executable binaries are ready immediately upon load. This removes the runtime PTX dependency, improves backwards compatibility, and reduces load times. PiperOrigin-RevId: 958260793
| Commit: | fa9dc26 | |
|---|---|---|
| Author: | Eusebio Durán Montaña | |
| Committer: | jax authors | |
Implement JAX-Triton Ahead-of-Time (AOT) Compilation & Serialization PiperOrigin-RevId: 958260793
| Commit: | 2050b2d | |
|---|---|---|
| Author: | Eusebio Durán Montaña | |
| Committer: | jax authors | |
Implement JAX-Triton Ahead-of-Time (AOT) Compilation & Serialization PiperOrigin-RevId: 957098418
| Commit: | 4d786de | |
|---|---|---|
| Author: | Allan Renucci | |
| Committer: | jax authors | |
[Mosaic GPU] Move GPU module loading and function resolution out of runtime. This change removes the required compilation of the initialization function for GPU kernels and loads modules directly in custom call handling. * Extract GPU binary and kernel metadata (name, dynamic shared memory size, and cluster size) as module attributes during launch lowering instead of generating a compiled initialization function. * Serialize the GPU binary and metadata directly in MosaicGpuKernelProto. * Invoke CUDA Driver and NVSHMEM APIs directly from custom_call.cc during kernel initialization. * Remove mosaic_gpu_module_load and mosaic_gpu_get_function from runtime.cc. We bump the mosaic_gpu proto version which will trigger XLA recompilation if there is a version mismatch. PiperOrigin-RevId: 955981262
| Commit: | 3fb1a1e | |
|---|---|---|
| Author: | Allan Renucci | |
| Committer: | jax authors | |
[Mosaic GPU] Move GPU module loading and function resolution out of JIT runtime. This change removes the required JIT compilation of the initialization function for GPU kernels and loads modules directly in custom call handling. * Extract GPU binary and kernel metadata (name, dynamic shared memory size, and cluster size) as module attributes during launch lowering instead of generating a JIT-compiled initialization function. * Serialize the GPU binary and metadata directly in MosaicGpuKernelProto. * Invoke CUDA Driver and NVSHMEM APIs directly from custom_call.cc during kernel initialization without resolving runtime library symbols. * Remove mosaic_gpu_module_load and mosaic_gpu_get_function from runtime.cc. PiperOrigin-RevId: 952007156
| Commit: | 6d68026 | |
|---|---|---|
| Author: | Olli Lupton | |
| Committer: | Olli Lupton | |
Mosaic GPU: pre-launch check for multicast support This gives a much better error message from kernels using multimem PTX instructions on incompatible systems.
| Commit: | 9fa9c3b | |
|---|---|---|
| Author: | Levon Ter-Grigoryan | |
| Committer: | jax authors | |
[Mosaic:GPU] Rename is_comm_used as is_nvshmem_used PiperOrigin-RevId: 878508055
| Commit: | a3ffef5 | |
|---|---|---|
| Author: | Aliia Khasanova | |
| Committer: | jax authors | |
Add serialization support for Mosaic GPU kernels. This change introduces a proto definition to serialize the state of a compiled Mosaic GPU kernel, including the object file and function names. The `CustomCallResources` struct is updated to stroe the kernel hash and include serialization and deserialization methods, allowing it to be used with XLA FFI's type registry. Also introduces `GetOrCreateKernel` to deduplicate kernels based on their kernel_hash (preserved during de/serialization). PiperOrigin-RevId: 877392204
| Commit: | 91223b9 | |
|---|---|---|
| Author: | Ilya Tikhonovskiy | |
| Committer: | jax authors | |
Follow up pr after triton integration cl and tokamax pr. Remove obsolete cluster dimensions from TritonKernel proto. PiperOrigin-RevId: 856837148
| Commit: | 81fe5dd | |
|---|---|---|
| Author: | Ilya Tikhonovskiy | |
| Committer: | jax authors | |
Integrate Triton up to 8d445186 https://github.com/openxla/triton/tree/triton_integrate_branch-1.15 PiperOrigin-RevId: 843206657
| Commit: | 72357ce | |
|---|---|---|
| Author: | Peter Hawkins | |
| Committer: | jax authors | |
[JAX] Migrate most remaining C++ code in jax to the jax:: namespace. PiperOrigin-RevId: 790476584
The documentation is generated from this commit.
| Commit: | dd1eac8 | |
|---|---|---|
| Author: | Peter Hawkins | |
| Committer: | jax authors | |
[JAX] Migrate most remaining C++ code in jax to the jax:: namespace. PiperOrigin-RevId: 789738942
The documentation is generated from this commit.
| Commit: | 39d7c38 | |
|---|---|---|
| Author: | Peter Hawkins | |
| Committer: | jax authors | |
Move contents of jaxlib/xla into jaxlib/ Having the directory structure of the jaxlib wheel be different to the source tree confuses type checkers such as mypy, since sometimes they find type stubs in the installed jaxlib wheel, and sometimes from the installed source tree. Instead: * don't include type stubs in the jaxlib wheel * don't install the jaxlib wheel as part of pre-commit * make sure that the location of type stubs (and the underlying libraries) is in the same position in the `jaxlib/` directory of the JAX source tree as it would be for the jaxlib wheel when installed. For now, we leave some stubs that forward from the old locations to the new locations for certain headers and modules. These will be removed after migrating some users. PiperOrigin-RevId: 750650528
| Commit: | f02a204 | |
|---|---|---|
| Author: | Peter Hawkins | |
| Committer: | jax authors | |
Move contents of jaxlib/xla into jaxlib/ Having the directory structure of the jaxlib wheel be different to the source tree is confusing type checkers such as mypy. PiperOrigin-RevId: 750398365
| Commit: | 588b693 | |
|---|---|---|
| Author: | Peter Hawkins | |
| Committer: | jax authors | |
[JAX] [XLA:Python] Migrate more Python modules to JAX. PiperOrigin-RevId: 740507886
| Commit: | 51560bf | |
|---|---|---|
| Author: | Peter Hawkins | |
| Committer: | jax authors | |
[JAX] [XLA:Python] Migrate pytree module to JAX. PiperOrigin-RevId: 740142231
| Commit: | ab3c1b5 | |
|---|---|---|
| Author: | jax authors | |
| Committer: | jax authors | |
[triton] Pass cluster_dims to TritonKernel and use cuLaunchKernel if size <= 1 PiperOrigin-RevId: 599809560
| Commit: | 3bb41c3 | |
|---|---|---|
| Author: | Peter Hawkins | |
| Committer: | jax authors | |
Move jaxlib/ under jax/_src/lib. Prototype only, not ready for review. PiperOrigin-RevId: 570062727
| Commit: | 4ac2bdc | |
|---|---|---|
| Author: | Chris Jones | |
| Committer: | jax authors | |
[jax_triton] Add user-specified `name` field to serialized format. PiperOrigin-RevId: 557415723
| Commit: | 714156d | |
|---|---|---|
| Author: | Chris Jones | |
| Committer: | jax authors | |
[jax_triton] Add support for float scalar inputs. Python `float`s are inferred as "f64". Values can be passed as "f32" using `np.float32(value)`. PiperOrigin-RevId: 552036612
| Commit: | f238667 | |
|---|---|---|
| Author: | Chris Jones | |
| Committer: | jax authors | |
Make JAX-Triton calls serializable. PiperOrigin-RevId: 542524794