Tzu-Wei Sung
5179642eb5
[Mosaic] Rename dep name.
...
PiperOrigin-RevId: 732985217
2025-03-03 11:01:25 -08:00
William S. Moses
8262987a1c
Fix build dependencies
...
PiperOrigin-RevId: 731330542
2025-02-26 08:38:31 -08:00
jax authors
b510127a13
Internal compatibility change
...
PiperOrigin-RevId: 729428478
2025-02-21 01:21:56 -08:00
jax authors
eaceac3bf9
[Pallas] Reductions with replicated axes.
...
PiperOrigin-RevId: 727292293
2025-02-15 07:41:16 -08:00
Tzu-Wei Sung
d4758b6d5e
[Mosaic][NFC] Factor out xla-array related utils in a separate file.
...
Also added tests.
PiperOrigin-RevId: 721424194
2025-01-30 09:49:41 -08:00
Sergei Lebedev
4221f109d1
[mosaic] Extracted serialization pass traversal logic into a reusable function
...
I will use it to implement Mosaic GPU serialization pass in a follow up.
PiperOrigin-RevId: 716156650
2025-01-16 02:58:06 -08:00
Tzu-Wei Sung
bf94389b08
[Mosaic] Use tpu::CreateMask for getX32VmaskByPaddingEnd.
...
It was cmp + iota before.
PiperOrigin-RevId: 713240888
2025-01-08 04:18:53 -08:00
Tzu-Wei Sung
57b21541a2
[Mosaic] NFC: Pull out vreg related functions to util.
...
These functions are related to vreg manipulation and are used in different rules.
PiperOrigin-RevId: 711484002
2025-01-02 11:50:19 -08:00
Sergei Lebedev
8987867faa
[mosaic_gpu] Include Mosaic GPU dialect fiels into jaxlib
2024-12-23 13:46:25 +00:00
Sergei Lebedev
a14e6968bf
[mosaic] Migrated the serialization pass from codegen to pass_boilerplate.h
...
This prepares teh generalization of the serialization pass to handle both
Mosaic TPU and GPU.
PiperOrigin-RevId: 705628923
2024-12-12 14:19:36 -08:00
Naums Mogers
0d36b0b433
[Mosaic] Add target core type parameter to tpu.sem_signal
...
Adds the optional core type parameter to `tpu.sem_signal` for cross-core signalling.
If the target core type is not provided, the target core type is assumed to be that of the core issuing the signal.
The issuing core type is determined based on the core type annotation of the parent function; if the annotation is not provided, the issuing core type is assumed to be TensorCore.
PiperOrigin-RevId: 698129842
2024-11-19 13:40:13 -08:00
Praveen Batra
8296f6e0ba
[Mosaic] Add extension files for infer/apply vector layout.
...
PiperOrigin-RevId: 691868278
2024-10-31 11:08:37 -07:00
Peter Hawkins
6d1f51e63d
Clean up BUILD files.
...
PiperOrigin-RevId: 667604964
2024-08-26 09:11:17 -07:00
Jevin Jiang
aa16485457
[XLA:Mosaic] Support memref shapecast.
...
This cl supports memref shapecast:
1. if tile is (1, 128), we support shapecast on any dim.
2. if shapecast on sublane dim, we only support tile aligned shape.
3. if shapecast on non-tiling dim, we support any shapecast.
4. all other cases would be considered as invalid memref shapecast.
PiperOrigin-RevId: 651924552
2024-07-12 17:05:03 -07:00
Kyle Lucke
84d748f43c
Stop using xla/statusor.h now that it just contains an alias for absl::Status.
...
In some situations, this meant also changing unrelated files to directly include tsl/platform/statusor.h to get the definitions for TF_ASSIGN_OR_RETURN, etc., where they were getting transitively included for free.
PiperOrigin-RevId: 645169743
2024-06-20 15:09:40 -07:00
jax authors
92d892b425
Adds rewrite patterns for arith
and math
operations with bf16
operands/results that are not supported by the underlying hardware.
...
PiperOrigin-RevId: 635865752
2024-05-21 11:08:19 -07:00
jax authors
8c2425e571
Adds rewrite patterns to LinalgVectorizationPass to eliminate transfer_read and transfer_write ops.
...
PiperOrigin-RevId: 628500668
2024-04-26 13:51:04 -07:00
Tomás Longeri
027c24e602
[Mosaic] Remove Python implementation of apply_vector_layout and infer_memref_layout.
...
PiperOrigin-RevId: 597332393
2024-01-10 13:00:21 -08:00
Blake Hechtman
1cba270f14
[XLA:MOSAIC] implement downcast from s32 to s8 correctly
...
PiperOrigin-RevId: 589830744
2023-12-11 08:20:53 -08:00
Tomás Longeri
250486f13e
[Mosaic] Expose C API for apply-vector-layout's assemble
, disassemble
, relayout
and applyLayoutOp
...
PiperOrigin-RevId: 580405900
2023-11-07 22:32:02 -08:00
Tomás Longeri
1c1dd7c8c7
[Mosaic] Expose C API for VectorLayout, VRegDataBounds
...
This is in preparation for Python bindings
PiperOrigin-RevId: 579355000
2023-11-03 18:24:16 -07:00
Tomás Longeri
a3a5926225
[Mosaic] apply_vector_layout rewrite (6): arith.constant
...
PiperOrigin-RevId: 567407089
2023-09-21 13:54:16 -07:00
Tomás Longeri
838f59e576
[Mosaic] apply_vector_layout C++ rewrite (4) Elementwise ops
...
PiperOrigin-RevId: 565255860
2023-09-13 22:19:10 -07:00
Tomás Longeri
a6eed40f24
[MOSAIC] apply_vector_layout C++ rewrite (3) applyLayoutOp and relayout
...
PiperOrigin-RevId: 563556815
2023-09-07 15:08:30 -07:00
Tomás Longeri
d02b59e410
[MOSAIC] apply_vector_layout C++ rewrite (1) VectorLayout functions
...
PiperOrigin-RevId: 561237760
2023-08-29 22:57:13 -07:00
Jevin Jiang
046bcc0ad9
[Mosaic] Add missing headers in linalg vectorization.
...
PiperOrigin-RevId: 560865251
2023-08-28 17:39:47 -07:00
Richard Levasseur
f891cbf64b
Load Python rules from rules_python
...
PiperOrigin-RevId: 559789250
2023-08-24 10:22:57 -07:00
Sharad Vikram
3baa6e7a89
Enable building jaxlib w/ Mosaic
...
PiperOrigin-RevId: 551159246
2023-07-26 03:59:30 -07:00
Sharad Vikram
3d556b7a19
Add Mosaic to Jaxlib and expose bindings in jax.experimental.mosaic
...
PiperOrigin-RevId: 549801858
2023-07-20 18:28:51 -07:00