
    uki}             	         d dl mZ d dlmZ d dlZd dlmZ d dlZd dlZd dl	m
Z
mZmZ d dlZd dlZd dlmZ d dlmZ d dlmZ d d	lmZ d d
lmZ d dlmZ d dlmZmZmZ d dlmZmZ d dlmZ d dlm Z  d dlm!Z! d dlm"Z" d dl#m$Z$ d dl#m%Z% d dl#m&Z' d dl(m)Z)m*Z*m+Z+ d dl,m-Z- d dl,m.Z. d dl,m/Z/ d dl,m0Z0 d dl,m1Z1 d dl2m3Z3 d dl4m5Z5 d dl4m6Z6 d dl7m8Z9 d dl:m;Z;m<Z< d  Z= e=e.        e=e/        e=e0        e=e1       d!d"dd#Z>dd$Z? G d% d&ej                        ZAd!d!ddd'	 	 	 	 	 	 	 	 	 	 	 dd(ZB G d) d*ej                        ZCd!d!d!ddd+	 	 	 	 	 	 	 	 	 	 	 	 	 dd,ZDdd-ZEdd.ZFdd/ZGdd0ZHed!dd1	 	 	 dd2       ZIed!dd1	 	 	 dd3       ZIed4d!dd5	 	 	 dd6       ZId4d!dd5	 	 	 dd7ZId!d4dd8	 	 	 	 	 	 	 	 	 dd9ZJ G d: d;ej                        ZKed!ddd<	 	 	 	 	 	 	 	 	 	 	 dd=       ZLed!ddd<	 	 	 	 	 	 	 	 	 	 	 dd>       ZLed!d!ddd?	 	 	 	 	 	 	 	 	 	 	 dd@       ZLd!d!ddd?	 	 	 	 	 	 	 	 	 	 	 ddAZLdBdCd4dD	 	 	 	 	 	 	 	 	 ddEZMd4d4d4d4d4dF	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddGZNd!dH	 	 	 	 	 ddIZOddJZPdKdLdMdNZQ	 ddOZRdP ZSdQ ZTdR ZU	 	 ddSZV eeVe%j                        ZXdT ZYdU ZZdV Z[dW Z\dX Z] eXe)e*z  fdYeYdZ      Z^eZe j                  e^<    e"j                  e^e[        e"j                  e^e\dK[        eRe^e]d\]       d^ Zad_ Zbd` Zc eXe)e)fdaeadbd4c      Zd e"j                  ed eecdL      dd[        e"j                  ed e"j                  ebd4e             df Zfdg Zgdh Zhdi Zidj Zjdk Zkdl Zl eVefe)e*z  fdYegdmd!e      Zmele j                  em<    e"j                  emeidK[        eRemekd\       dn Zndo Zo	 	 ddpZpdq Zq eVeoe)e*z  fdYendrd!e      Zreqe j                  er<    eRerep       ds Zsdt Ztdu Zu eVete)e*z  fdYesdvd!e      Zv e"j                  eveudK[       dw Zwdx Zx	 	 ddyZy eXe)e*z  e)e*z  fdaewdz      Zz e"j                  ezex        eRezey       d{ Z{d d|Z|d} Z}d~ Z~d Zd Zd ZddZd Z eVee)e*z  fdYe~dd!e      Zee j                  e<    e"j                  e e"j                  e}d!e              e"j                  eed[        eRee       	 d	 	 	 ddZddZ ej                  d      dd       Zd Zd Zd Zd Zd Z eXej                  hfded      Z e"j                  e e"j                  ed4e              eReed\       ddZd Zd Zd ZddZ eVee)e*z  fdYedd!e      Z e"j                  ee        eRee       dd	 	 	 ddZd Zd Zd Z eVee)e*z  e+fdaedd!d4      Z eRee       d Zd Zd Zd Z eVee)e*z  fdYedd!e      Zee j                  e<    e"j                  e e"j                  e             d Zd Zd Z eVee)e*z  fdYedd!e      Z e"j                  eedK[       d Zd Z ejL                  d      ddd       Zd Zd Zdd	 ddZd Z eVee)e*z  fdYedd!e      Zee j                  e<    eRee       d Zd Zd Z eXe)e)fded      Z e"j                  e eedL      dd[        e"j                  e e"j                  ed4e             d4ddZd Zd Zd Zd Zd Z ejn                  ejp                         ejn                  ejr                         ejn                  ejt                         ejn                  ejv                        hZd Z eVee)e*z  e)e*z  fded      Z e j~                  eed        ee j                  e<   ee!j                  e<    e"j                  ee        e"j                  eedK[       d Zd ZdÄ Z eVee)e*z  fdYedd!e      Z eReeī       dń ZdƄ ZdǄ ZdȄ ZdɄ Zdʄ Zd˄ Zd̄ Zd̈́ Z eXe)e*z  e)e*z  e)e*z  e)e*z  fdedϫ      Zee j                  e<   ee j                  e<   ee!j                  e<    e"j                  eedK[        e"j                  e eedLЫ      dd[        e"j                  e eedMЫ      dѬ[        e"j                  e e"j                  ed4e             dd҄ZАddӄZѐddԄZҐddՄZӐddքZԐddׄZՐddd؄Z֐dddلZאd	dڄZؐd	dۄZِd
d܄Zڐdd݄ZېddބZd߄ Zd Zd Zd Zd ZddZddZddZd Zd Z	 	 	 ddZy(      )annotations)CallableN)partial)AnyLiteraloverload)ad_util)api)config)core)dispatch)dtypes)ShapedArrayis_constant_dimis_constant_shape)sdy_sharding_rule_to_mlirstr_to_sdy_sharding_rule)ffi)ad)batching)mlir)control_flow)lax)utils)_float_complex_int)cuda_versions)
gpu_linalg)
gpu_solver)
gpu_sparse)lapack)ir)chlo)hlo)PartitionSpec)Array	ArrayLikec                   t        | d      rJ| j                         j                         D ])  \  }}|D ]  \  }}}t        j                  ||||       ! + t        | d      r+| j                         D ]  }t        j                  |        y y )Nregistrations)platformapi_versionbatch_partitionable_targets)hasattrr*   itemsr   register_ffi_targetr-   *register_ffi_target_as_batch_partitionable)moduler+   targetsnamevaluer,   s         N/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/_src/lax/linalg.pyregister_module_custom_callsr7   8   s    V_%#11399; 
'&- 

"${%(	



 V23224 ;	44T:; 4    T)symmetrize_inputc               X    |rt        |       } t        t        j                  |             S )a  Cholesky decomposition.

  Computes the Cholesky decomposition

  .. math::
    A = L . L^H

  of square matrices, :math:`A`, such that :math:`L`
  is lower triangular. The matrices of :math:`A` must be positive-definite and
  either Hermitian, if complex, or symmetric, if real.

  Args:
    x: A batch of square Hermitian (symmetric if real) positive-definite
      matrices with shape ``[..., n, n]``.
    symmetrize_input: If ``True``, the matrix is symmetrized before Cholesky
      decomposition by computing :math:`\frac{1}{2}(x + x^H)`. If ``False``,
      only the lower triangle of ``x`` is used; the upper triangle is ignored
      and not accessed.

  Returns:
    The Cholesky decomposition as a matrix with the same dtype as ``x`` and
    shape ``[..., n, n]``. If Cholesky decomposition fails, returns a matrix
    full of NaNs. The behavior on failure may change in the future.
  )
symmetrize_tril
cholesky_pbind)xr9   s     r6   choleskyr@   L   s$    2 1A	zq!	""r8   c                `    t        j                  | |      \  } }t        j                  | |      S )a  Cholesky rank-1 update.

  Given a Cholesky decomposition :math:`A = R.T \, R` and a vector :math:`w`,
  computes the Cholesky decomposition of :math:`A + w \, w.T` in :math:`O(N^2)`
  time.

  Args:
    r_matrix: An upper-triangular matrix (R) such that :math:`A = R^T \, R`.
    w_vector: A vector :math:`w` for rank-1 update.

  Returns:
    A new upper-triangular matrix :math:`R` defining the Cholesky decomposition
    of :math:`A + w \, w^T`.
  )r   standard_insert_pvarycholesky_update_pr>   )r_matrixw_vectors     r6   cholesky_updaterF   j   s.     11(HE(H			(	33r8   c                      e Zd ZdZdZdZdZy)EigImplementationz&Enum for eigendecomposition algorithm.cusolvermagmar"   N)__name__
__module____qualname____doc__CUSOLVERMAGMALAPACK r8   r6   rH   rH   }   s    .(
%&r8   rH   )compute_left_eigenvectorscompute_right_eigenvectorsimplementation	use_magmac                   |>t        j                  dt        d       |rt        j                  nt        j
                  }t        j                  | |||      S )a  Eigendecomposition of a general matrix.

  Nonsymmetric eigendecomposition is only implemented on CPU and GPU. On GPU,
  the default implementation calls LAPACK directly on the host CPU, but an
  experimental GPU implementation using `MAGMA <https://icl.utk.edu/magma/>`_
  is also available. The MAGMA implementation is typically slower than the
  equivalent LAPACK implementation for small matrices (less than about 2048),
  but it may perform better for larger matrices.

  To enable the MAGMA implementation, you must install MAGMA yourself (there
  are Debian and conda-forge packages, or you can build from source). Then set
  the ``use_magma`` argument to ``True``, or set the ``jax_use_magma``
  configuration variable to ``"on"`` or ``"auto"``:

  .. code-block:: python

      jax.config.update('jax_use_magma', 'on')

  JAX will try to ``dlopen`` the installed MAGMA shared library, raising an
  error if it is not found. To explicitly specify the path to the MAGMA
  library, set the environment variable `JAX_GPU_MAGMA_PATH` to the full
  installation path.

  If ``jax_use_magma`` is set to ``"auto"``, the MAGMA implementation will
  be used if the library can be found, and the input matrix is sufficiently
  large (>= 2048x2048).

  Args:
    x: A batch of square matrices with shape ``[..., n, n]``.
    compute_left_eigenvectors: If true, the left eigenvectors will be computed.
    compute_right_eigenvectors: If true, the right eigenvectors will be
      computed.
    use_magma: Deprecated, please use ``implementation`` instead. Locally
      override the ``jax_use_magma`` flag. If ``True``, the eigendecomposition
      is computed using MAGMA. If ``False``, the computation is done using
      LAPACK on to the host CPU. If ``None`` (default), the behavior is
      controlled by the ``jax_use_magma`` flag. This argument is only used on
      GPU. Will be removed in JAX 0.9.
    implementation: Controls the choice of eigendecomposition algorithm. If
    ``LAPACK``, the computation will be performed using LAPACK on the host CPU.
      If ``MAGMA``, the computation will be performed using the MAGMA library on
      the GPU. If ``CUSOLVER``, the computation will be performed using the
      Cusolver library on the GPU. The ``CUSOLVER`` implementation requires
      Cusolver 11.7.1 (from CUDA 12.6 update 2) to be installed, and does not
      support computing left eigenvectors.
      If ``None`` (default), an automatic choice will be made, depending on the
      Cusolver version, whether left eigenvectors were requested, and the
      ``jax_use_magma`` configuration variable.

  Returns:
    The eigendecomposition of ``x``, which is a tuple of the form
    ``(w, vl, vr)`` where ``w`` are the eigenvalues, ``vl`` are the left
    eigenvectors, and ``vr`` are the right eigenvectors. ``vl`` and ``vr`` are
    optional and will only be included if ``compute_left_eigenvectors`` or
    ``compute_right_eigenvectors`` respectively are ``True``.

    If the eigendecomposition fails, then arrays full of NaNs will be returned
    for that batch element.
  zSuse_magma is deprecated, please use implementation=EigImplementation.MAGMA instead.   )
stacklevel)rS   rT   rU   )warningswarnDeprecationWarningrH   rP   rQ   eig_pr>   )r?   rS   rT   rU   rV   s        r6   eigr^      sa    F MM	;	 $-2C2J2J  
A1J/I#1 
 
3 3r8   c                      e Zd ZdZdZdZdZy)EighImplementationz:Implementation for symmetric/Hermitian eigendecomposition.qrjacobiqdwhN)rK   rL   rM   rN   QRJACOBIQDWHrR   r8   r6   r`   r`      s    B"&	$r8   r`   )lowerr9   sort_eigenvaluessubset_by_indexrU   c               ^    |rt        |       } t        j                  | ||||      \  }}||fS )af  Eigendecomposition of a Hermitian matrix.

  Computes the eigenvectors and eigenvalues of a complex Hermitian or real
  symmetric square matrix.

  Args:
    x: A batch of square complex Hermitian or real symmetric matrices with shape
      ``[..., n, n]``.
    lower: If ``symmetrize_input`` is ``False``, describes which triangle of the
      input matrix to use. If ``symmetrize_input`` is ``False``, only the
      triangle given by ``lower`` is accessed; the other triangle is ignored and
      not accessed.
    symmetrize_input: If ``True``, the matrix is symmetrized before the
      eigendecomposition by computing :math:`\frac{1}{2}(x + x^H)`.
    sort_eigenvalues: If ``True``, the eigenvalues will be sorted in ascending
      order. If ``False`` the eigenvalues are returned in an
      implementation-defined order.
    subset_by_index: Optional 2-tuple [start, end] indicating the range of
      indices of eigenvalues to compute. For example, is ``range_select`` =
      [n-2,n], then ``eigh`` computes the two largest eigenvalues and their
      eigenvectors.
    implementation: Optional implementation selection. ``QR`` uses QR-based
      decomposition (default for CPU/GPU). ``JACOBI`` uses Jacobi iteration
      (GPU/TPU only). ``QDWH`` uses QDWH spectral divide-and-conquer
      (default on TPU, TPU only).

  Returns:
    A tuple ``(v, w)``.

    ``v`` is an array with the same dtype as ``x`` such that ``v[..., :, i]`` is
    the normalized eigenvector corresponding to eigenvalue ``w[..., i]``.

    ``w`` is an array with the same dtype as ``x`` (or its real counterpart if
    complex) with shape ``[..., d]`` containing the eigenvalues of ``x`` in
    ascending order(each repeated according to its multiplicity).
    If ``subset_by_index`` is ``None`` then ``d`` is equal to ``n``. Otherwise
    ``d`` is equal to ``subset_by_index[1] - subset_by_index[0]``.
  rg   rh   ri   	algorithm)r;   eigh_pr>   )r?   rg   r9   rh   ri   rU   vws           r6   eighrp      sB    ^ 1A	'% 
 
$!Q 
A+r8   c                ,    t         j                  |       S )a  Reduces a square matrix to upper Hessenberg form.

  Currently implemented on CPU only.

  Args:
    a: A floating point or complex square matrix or batch of matrices.

  Returns:
    A ``(a, taus)`` pair, where the upper triangle and first subdiagonal of
    ``a`` contain the upper Hessenberg matrix, and the elements below the first
    subdiagonal contain the Householder reflectors. For each Householder
    reflector ``taus`` contains the scalar factors of the elementary Householder
    reflectors.
  )hessenberg_pr>   )as    r6   
hessenbergrt     s     
		1	r8   c                `    t        j                  | |      \  } }t        j                  | |      S )a  Product of elementary Householder reflectors.

  Args:
    a: A matrix with shape ``[..., m, n]``, whose lower triangle contains
      elementary Householder reflectors.
    taus: A vector with shape ``[..., k]``, where ``k < min(m, n)``, containing
      the scalar factors of the elementary Householder reflectors.

  Returns:
    A batch of orthogonal (unitary) matrices with the same shape as ``a``,
    containing the products of the elementary Householder reflectors.
  )r   rB   householder_product_pr>   )rs   tauss     r6   householder_productrx   *  s-     &&q$/'!T		#	#At	,,r8   c                ,    t         j                  |       S )a  LU decomposition with partial pivoting.

  Computes the matrix decomposition:

  .. math::
    P \, A = L \, U

  where :math:`P` is a permutation of the rows of :math:`A`, :math:`L` is a
  lower-triangular matrix with unit-diagonal elements, and :math:`U` is an
  upper-triangular matrix.

  Args:
    x: A batch of matrices with shape ``[..., m, n]``.

  Returns:
    A tuple ``(lu, pivots, permutation)``.

    ``lu`` is a batch of matrices with the same shape and dtype as ``x``
    containing the :math:`L` matrix in its lower triangle and the :math:`U`
    matrix in its upper triangle. The (unit) diagonal elements of :math:`L` are
    not represented explicitly.

    ``pivots`` is an int32 array with shape ``[..., min(m, n)]`` representing a
    sequence of row swaps that should be performed on :math:`A`.

    ``permutation`` is an alternative representation of the sequence of row
    swaps as a permutation, represented as an int32 array with shape
    ``[..., m]``.
  )lu_pr>   r?   s    r6   lur|   ;  s    < 
1r8   c                0    t         j                  | |      S )a  Converts the pivots (row swaps) returned by LU to a permutation.

  We build a permutation rather than applying `pivots` directly to the rows
  of a matrix because lax loops aren't differentiable.

  Args:
    pivots: an int32 array of shape (..., k) of row swaps to perform
    permutation_size: the size of the output permutation. Has to be >= k.

  Returns:
    An int32 array of shape (..., permutation_size).
  )permutation_size)lu_pivots_to_permutation_pr>   )pivotsr~   s     r6   lu_pivots_to_permutationr   \  s"     
$	(	(/ 
) 
1 1r8   )full_matricesrV   c                    y NrR   r?   pivotingr   rV   s       r6   ra   ra   m       r8   c                    y r   rR   r   s       r6   ra   ra   r  r   r8   Fr   r   rV   c                    y r   rR   r   s       r6   ra   ra   w  s     r8   c               V    t         j                  | |||      ^}}}|r|||d   fS ||fS )a  QR decomposition.

  Computes the QR decomposition

  .. math::
    A = Q \, R

  of matrices :math:`A`, such that :math:`Q` is a unitary (orthogonal) matrix,
  and :math:`R` is an upper-triangular matrix.

  Args:
    x: A batch of matrices with shape ``[..., m, n]``.
    pivoting: Allows the QR decomposition to be rank-revealing. If ``True``,
      compute the column pivoted decomposition ``A[:, P] = Q @ R``, where ``P``
      is chosen such that the diagonal of ``R`` is non-increasing. Currently
      supported on CPU and GPU backends only.
    full_matrices: Determines if full or reduced matrices are returned; see
      below.
    use_magma: Locally override the ``jax_use_magma`` flag. If ``True``, the
      pivoted `qr` factorization is computed using MAGMA. If ``False``, the
      computation is done using LAPACK on the host CPU. If ``None`` (default),
      the behavior is controlled by the ``jax_use_magma`` flag. This argument is
      only used on GPU.

  Returns:
    A pair of arrays ``(q, r)``, if ``pivoting=False``, otherwise ``(q, r, p)``.

    Array ``q`` is a unitary (orthogonal) matrix,
    with shape ``[..., m, m]`` if ``full_matrices=True``, or
    ``[..., m, min(m, n)]`` if ``full_matrices=False``.

    Array ``r`` is an upper-triangular matrix with shape ``[..., m, n]`` if
    ``full_matrices=True``, or ``[..., min(m, n), n]`` if
    ``full_matrices=False``.

    Array ``p`` is an index vector with shape [..., n]

  Notes:
    - `MAGMA <https://icl.utk.edu/magma/>`_ support is experimental - see
      :func:`jax.lax.linalg.eig` for further assumptions and limitations.
    - If ``jax_use_magma`` is set to ``"auto"``, the MAGMA implementation will
      be used if the library can be found, and the input matrix is sufficiently
      large (has at least 2048 columns).
  r   r   )qr_pr>   )r?   r   r   rV   qrps          r6   ra   ra   }  sC    ^ YYq8=!*  ,(!Qa1:	
A+r8   compute_schur_vectorssort_eig_valsselect_callablec               4    t         j                  | |||      S )a  Schur decomposition.

  Only implemented on CPU.

  Computes the Schur decomposition:

  .. math::
    A = Q \, U \, Q^{-H}

  for a square matrix :math:`A`.

  Args:
    x: A batch of square matrices with shape ``[..., m, m]``.
    compute_schur_vectors: If ``True``, compute the Schur vectors ::math:`Q`,
      otherwise only :math:`U` is computed.
    sort_eig_vals: Unused.
    select_callable: Unused.

  Returns:
    A pair of arrays ``U, Q``, if ``compute_schur_vectors=True``, otherwise
    only ``U`` is returned.
  r   )schur_pr>   )r?   r   r   r   s       r6   schurr     s&    : 
1!%	 
 
' 'r8   c                       e Zd ZdZdZdZdZdZy)SvdAlgorithmzEnum for SVD algorithm.defaultrd   JacobipolarN)rK   rL   rM   rN   DEFAULTrd   re   POLARrR   r8   r6   r   r     s    '"&
%r8   r   )r   ri   rl   c                    y r   rR   r?   r   
compute_uvri   rl   s        r6   svdr          r8   c                    y r   rR   r   s        r6   r   r     r   r8   r   r   ri   rl   c                    y r   rR   r   s        r6   r   r     r   r8   c               \    t         j                  | ||||      }|r|\  }}}|||fS |\  }|S )a  Singular value decomposition.

  Computes the singular value decomposition of an input matrix.

  Args:
    x: A batch of matrices with shape ``[..., m, n]``.
    full_matrices: Determines if full or reduced matrices are returned.
    compute_uv: If ``True``, returns the left singular vectors, the singular
      values and the adjoint of the right singular vectors. Otherwise, only
      the singular values are returned.
    subset_by_index: If ``None``, the entire matrix is returned. Otherwise,
      returns the singular values and vectors for the given range of indices.
    algorithm: The SVD algorithm to use. Must be ``None`` or a value from
      :class:`~jax.lax.linalg.SvdAlgorithm`.

  Returns:
    The singular values if ``compute_uv`` is ``False``, otherwise returns a
    triple containing the left singular vectors, the singular values, and the
    adjoint of the right singular vectors.
  r   )svd_pr>   )	r?   r   r   ri   rl   resultsurn   s	            r6   r   r     sM    8 ::!%  & GAq!a7N	BAHr8         ?        )alphabetasymmetrize_outputc               >   t        j                  | |      \  } }t        j                  | |||      }|rht	        j
                  t        |d      g t        |j                  dz
        |j                  dz
  |j                  dz
        }t        |d      |z   }|S )a  Symmetric product.

  Computes the symmetric product

  .. math::
    \alpha \, A \, A^T + \beta \, C

  where :math:`A` is a rectangular matrix and :math:`C` is a symmetric matrix.

  Args:
    a_matrix: A batch of matrices with shape ``[..., m, n]``.
    c_matrix: A batch of matrices with shape ``[..., m, m]``.
    alpha: A scalar.
    beta: A scalar.
    symmetrize_output: If ``True``, the upper triangle of the output is
      replaced with its transpose.

  Returns:
    A batch of matrices with shape ``[..., m, m]`` where only the lower
    triangle is guaranteed to include the correct values on all platforms. If
    ``symmetrize_output`` is ``True``, the upper triangle is filled with the
    transpose of the lower triangle, and the whole matrix is valid.
  )r   r   krX      r   )	r   rB   symmetric_product_pr>   r   	transposer<   rangendim)a_matrixc_matrixr   r   r   r   
upper_halfs          r6   symmetric_productr   /  s    > 11(HE(H##Hhe$#O&fC%a
 C&++/C6;;?CEJ 6Q*,F	-r8   	left_siderg   transpose_aconjugate_aunit_diagonalc          	        |xr7 t        j                  t        j                  |       t        j
                        }t	        j                  |      t	        j                  |       dz
  k(  }|rt        j                  ||rdndf      }t        j                  | |      \  } }t        j                  | ||||||      }|r|r|d   n	|ddddf   }|S )	a  Triangular solve.

  Solves either the matrix equation

  .. math::
    \mathit{op}(A) . X = B

  if ``left_side`` is ``True`` or

  .. math::
    X . \mathit{op}(A) = B

  if ``left_side`` is ``False``.

  ``A`` must be a lower or upper triangular square matrix, and where
  :math:`\mathit{op}(A)` may either transpose :math:`A` if ``transpose_a``
  is ``True`` and/or take its complex conjugate if ``conjugate_a`` is ``True``.

  Args:
    a: A batch of matrices with shape ``[..., m, m]``.
    b: A batch of matrices with shape ``[..., m, n]`` if ``left_side`` is
      ``True`` or shape ``[..., n, m]`` otherwise.
    left_side: describes which of the two matrix equations to solve; see above.
    lower: describes which triangle of ``a`` should be used. The other triangle
      is ignored.
    transpose_a: if ``True``, the value of ``a`` is transposed.
    conjugate_a: if ``True``, the complex conjugate of ``a`` is used in the
      solve. Has no effect if ``a`` is real.
    unit_diagonal: if ``True``, the diagonal of ``a`` is assumed to be unit
      (all 1s) and not accessed.

  Returns:
    A batch of matrices the same shape and dtype as ``b``.
  r   r   r   .r   .r   N)r   
issubdtyper   dtypenpcomplexfloatingr   expand_dimsr   rB   triangular_solve_pr>   )	rs   br   rg   r   r   r   	singletonouts	            r6   triangular_solver   X  s    X S 1 1#))A,@R@R S+ggajBGGAJN*))B56A		#	#Aq	)$!Qiu+] 	  	<# "#f+CAIC	*r8   rg   c               V    t         j                  t        j                  |       |      S )aE  Reduces a symmetric/Hermitian matrix to tridiagonal form.

  Currently implemented on CPU and GPU only.

  Args:
    a: A floating point or complex matrix or batch of matrices.
    lower: Describes which triangle of the input matrices to use.
      The other triangle is ignored and not accessed.

  Returns:
    A ``(a, d, e, taus)`` tuple. If ``lower=True``, the diagonal and first
    subdiagonal of matrix (or batch of matrices) ``a`` contain the tridiagonal
    representation, and elements below the first subdiagonal contain the
    elementary Householder reflectors, where additionally ``d`` contains the
    diagonal of the matrix and ``e`` contains the first subdiagonal. If
    ``lower=False`` the diagonal and first superdiagonal of the matrix contains
    the tridiagonal representation, and elements above the first superdiagonal
    contain the elementary Householder reflectors, where additionally ``d``
    contains the diagonal of the matrix and ``e`` contains the first
    superdiagonal. ``taus`` contains the scalar factors of the elementary
    Householder reflectors.
  r   )tridiagonal_pr>   r   asarray)rs   rg   s     r6   tridiagonalr     s"    2 
		CKKN%		88r8   c                l    t        j                  | |||      \  } }}}t        j                  | |||      S )a  Computes the solution of a tridiagonal linear system.

  This function computes the solution of a tridiagonal linear system:

  .. math::
    A \, X = B

  Args:

    dl: A batch of vectors with shape ``[..., m]``.
      The lower diagonal of A: ``dl[i] := A[i, i-1]`` for i in ``[0,m)``.
      Note that ``dl[0] = 0``.
    d: A batch of vectors with shape ``[..., m]``.
      The middle diagonal of A: ``d[i]  := A[i, i]`` for i in ``[0,m)``.
    du: A batch of vectors with shape ``[..., m]``.
      The upper diagonal of A: ``du[i] := A[i, i+1]`` for i in ``[0,m)``.
      Note that ``dl[m - 1] = 0``.
    b: Right hand side matrix.

  Returns:
    Solution ``X`` of tridiagonal system.
  )r   rB   tridiagonal_solve_pr>   )dlddur   s       r6   tridiagonal_solver     s9    . ++B2q9,"aQ		!	!"aQ	//r8   cpucuhipr   cudarocmc                j    |D ].  }t         |   }t        j                  | t        ||      |       0 y )Ntarget_name_prefixr+   )_platform_prefix_mapr   register_loweringr   )primlowering_rulesupported_platformsr+   prefixs        r6   register_cpu_gpu_loweringr     s;     & h!(+F&9r8   c                v   g g c}t        t        ||            D ]  \  }\  }	}
|
j                  }t        |      |	k  rt	        d| d| d|	 d|       |s%t        |      |	k7  rt	        d| d| d|	 d|       j                  |d t        |      |	z
          |j                  |t        |      |	z
  d          t        fdD              s)t	        d| dD cg c]  }t        |       c} d	      t        d
          ||i |}| rt        fd|D              S t        |      z   S c c}w )NInput  to z must have rank at least z, but got shape=z must have a rank of exactly c              3  R   K   | ]  }t        d          t        |      k(     yw)r   N)len).0r   
batch_dimss     r6   	<genexpr>z$linalg_shape_rule.<locals>.<genexpr>  s#     >aSA3q6)>s   $'All inputs to z8 must have the same number of batch dimensions, but got z! batch dimensions for the inputs.r   c              3  :   K   | ]  }t        |      z     y wr   )tuple)r   r   r   s     r6   r   z$linalg_shape_rule.<locals>.<genexpr>  s     41eAh&4s   )	enumeratezipshaper   	TypeErrorappendallr   )multiple_resultssupports_batchingranksresult_shaper4   avalskwargsdimsirankavalr   r   r   r   s                 @r6   linalg_shape_ruler    s   *d"3ue#45 +oa$JJE
5zD1#T$8 ?'  Ut!31#T$<TF C'  e.SZ$./0KKc%j4'()*+ 
>:>	>

 $./qCF/0 1	 
 Z]#*d%f%#4444c
"" 0s    D6c                    ||i |}g }t        t        ||            D ]  \  }\  }	}
|
j                  j                  }|d t	        |      |	z
   |t	        |      |	z
  d  c}t        d |D              st        j                  d| d| d| d      |j                          |d   t        fd|dd  D              rt        j                  d	| d
| d      |d   j                  }| rJ|D cg c]>  }|j                  t        t              dt	        |      t	              z
  z  z          @ c}S t	        |      t	              z
  }|j                  t        t              d|z  z          S c c}w )Nc              3  $   K   | ]  }|d u  
 y wr   rR   )r   r   s     r6   r   z'linalg_sharding_rule.<locals>.<genexpr>  s     ,QqDy,s   r   r   z4 must be unsharded on non-batch dimensions, but got .r   c              3  (   K   | ]	  }|k7    y wr   rR   )r   r   
batch_specs     r6   r   z'linalg_sharding_rule.<locals>.<genexpr>  s     2Qj2s   r   r   z, must have the same batch sharding, but got r   spec)r   r   shardingr  r   r   r   ShardingTypeErrorr   anyupdatePr   )r   
shape_ruler   r4   r   r  output_shapesbatch_specsr  r  r  r  	rest_specr  r   r   r  s                   @r6   linalg_sharding_ruler    s    e.v.-+"3ue#45 #oa$==D !2#d)d"23T#d)d:J:K5LJ	,),,""1#T$  F!  z"# 1~*2+ab/22

 
 
J-q	  1X(
 	  	j!Gs1vJ/G$HHJ 	 	
  }J/D??E*$5$$F H?IIs   4AE8c                d     ||i |}t        j                  |g| }| r|gt        |      z  S |S r   )r   standard_vma_ruler   )r   r  r4   r   r  r  out_vmas          r6   linalg_vma_ruler    s@    e.v.-""40%0'9s=)))Nr8   c                   t        t        j                  | |||      }t        t        |||||      }	|rt        t        ||	||      }
nd }
t        t
        ||	|      }t        j                  |      }||_        |j                  t        t        j                  |             |r=|j                  t        t        j                  ||	|t        j                  |
|             nW|j                  t        t        j                   ||	|t        j                  |
t        t        j"                  |      d d d 
             |r+t        t$        j&                  |      t$        j(                  |<   |S )N)require_same)r   r   naryop_dtype_ruler  r  r  r   	Primitiver   def_implr   apply_primitivedef_abstract_eval	lax_utils#standard_multi_result_abstract_eval_standard_weak_type_rulestandard_abstract_evalr  r   expand_dims_batcherprimitive_batchers)result_dtypeaccepted_dtypesr   r   r4   r   r   r  
dtype_ruler  sharding_rulevma_ruler   s                r6   linalg_primitiver.    sG    	\?D!* )+<eD* .
E4IM M_&6
DI(		$*$--00$78	==tJ	(J(Jx	)*
 	i..j*00-d,,d3T4GH (/$$d),H%	+r8   c                :    | d   | d   k7  rt        d|  d      | S )Nr   r   z9The input to cholesky must be a square matrix. Got shape r	  
ValueErrorr   s    r6   _cholesky_shape_ruler3  D  s2    
1Xq

CE7!LN N	,r8   c                   | \  }|\  }t        t        j                  |            }d }t        ||dddd      }t	        j
                  | |t        ||ddd            t        j                  j                        }||fS )Nc           
        t        |       }|t        j                  t        j                  | d      t        j                  | j
                  | j                  d   | j                  d   f      z   t        |j                  dz
              z  S )Nr   r   rX   )	r<   r   r   _const_eyer   r   r   r   )Xls     r6   phiz_cholesky_jvp_rule.<locals>.phiQ  sj    aAs

1a388AGGaggbk1772;-GHHaffqj  r8   FT)r   r   r   rg   r   r   rg   	precision)r<   r=   r>   r   r   batch_matmul	PrecisionHIGHEST)primalstangentsr?   	sigma_dotLr:  tmpL_dots           r6   _cholesky_jvp_rulerG  K  s    "!*)JOOA! 	ID%)	7#


1c"2%t#= >%%'% 
E/r8   c                n    ~ t        j                  |t        j                  j	                  d            gS )NTr   )r%   r@   r#   BoolAttrget)ctxr?   s     r6   _cholesky_loweringrL  _  s&    	
,,q 5
6	77r8   c                   | j                   \  }| j                  \  }|j                  d d }t        j                  d|j
                        }t        |t        j                        }t        |||gddi      } || |t        d            \  }}	t        j                  |	t        j                  | d|      dd      }
t        | ||
||      gS )	Nr   	potrf_ffir   	avals_outoperand_output_aliasesT)uploEQSIGNED)avals_inrP  r   r"   prepare_lapack_callr   r   r   int32_linalg_ffi_lowering_matrix_uplo_attrr   compare_hlofull_like_aval_replace_not_ok_with_nan)rK  operandoperand_avalout_avalr   target_name	info_avalruler   infooks              r6   _cholesky_cpu_loweringre  d  s    ,,-,mm)(!!#2&***;8J8JK+*bhh/)	klI5N67V
=$c7):4)@A,&$d11#q)Dd ""
"3
B
I	JJr8   c               L   | j                   \  }| j                  \  }|j                  d d }t        |t        j
                        }t        | d||gddi      } || |d      \  }}	t        j                  |	t        j                  | d|      dd      }
t        | ||
||      gS )	Nr   solver_potrf_ffir   rO  Tr   rS  rT  )rU  rP  r   r   r   rW  rX  r   rZ  r[  r\  )rK  r]  r   r^  r_  r   ra  rb  r   rc  rd  s              r6   _cholesky_gpu_loweringrh  r  s    ,,-,mm)(!!#2&**bhh/)	!3 44DE)5y(A67V
=$ c7$/,&$d11#q)Dd ""
"3
B
I	JJr8   )rX   r@   r   )r   r   )r   c                V    | d   | d   k7  s|d   | d   k7  rt        d|  d| d      | S )Nr   r   zqRank-1 update to Cholesky decomposition takes a square matrix and a vector of the same size as input. Got shapes  and z insteadr0  )r_shapew_shapes     r6   _cholesky_update_shape_rulerm    sM    QZ71:wqz!9
	>>EYe)8	  
.r8   c                    d }	 	 	 	 	 	 	 	 	 	 dd}|j                   d   }t        |      D ]M  } || ||f   ||         \  }} || |d d f   |||      \  }}| j                  |d d f   j                  |      } O | S )Nc                    d }t        j                  d| j                        t        j                  d| j                        ft        j                  |dk(  fd|| |      S )z:Get coefs for Givens rotation in a numerically stable way.c                    t        |       }t        |      }t        j                  ||kD  ||      }| |z  } ||z  }dt        j                  | dz  |dz  z         z  }| |z  | |z  fS )Nr   rX   )absr   selectsqrt)r?   yabs_xabs_ydenominatorrhs         r6   _drotg_nonzeroz?_cholesky_update_jax_fn.<locals>._drotg.<locals>._drotg_nonzero  sq    !fe!feJJuu}eU;k;a;asxxQa((bVaR"W_r8   r   r   r   r   c                    S r   rR   )r?   rt  one_and_zeros     r6   <lambda>z9_cholesky_update_jax_fn.<locals>._drotg.<locals>.<lambda>  s    \ r8   )r   arrayr   r   cond)r?   rt  ry  r|  s      @r6   _drotgz'_cholesky_update_jax_fn.<locals>._drotg  s\     	177#
177#L 	Q)>1aA Ar8   c                .    || z  ||z  z
  ||z  || z  z   fS r   rR   )first_vectorsecond_vectorc_coefs_coefs       r6   _drotz&_cholesky_update_jax_fn.<locals>._drot  s2     	 66,!668 8r8   r   )
r  r'   r  r'   r  floatr  r  returntuple[Array, Array])r   r   atset)	Rzr  r  nr   cr   row_ks	            r6   _cholesky_update_jax_fnr    s    A"88*/88"8':8 ggaj!8 a!AqD'1Q4 DAqQq!tWaA&HE1	QT
uA 
(r8   c                    t        j                  |  dddd      }|j                  |j                        } ||||      d d S )N_cholesky_update_ffir   r   r   r   rQ  rP  )r   ffi_loweringreplacerU  )r   rK  rD   rE   rb  sub_ctxs         r6   "_cholesky_update_gpu_lowering_ruler    sO    			/00DE561
?$KK#,,K/'	gx	*2A	..r8   )rX   r   rF   )r   r   r   c               D    t        j                  |       }|fd|z   |z   z  S Nr   )r   to_complex_dtype)a_dtyperS   rT   _r   s        r6   _eig_dtype_ruler    s-     
!
!'
*%
Q225OO	PPr8   c               Z    | d   | d   k7  rt        d|  d      ||z   }| d d f| f|z  z   S )Nr   r   z4The input to eig must be a square matrix. Got shape r	  r   r0  )r   rS   rT   r  counts        r6   _eig_shape_ruler    sT     1Xq

>ugQGI I
#&@
@%
*%E)	))r8   c                    t        | r(t        j                  j                  j                        S t        j                  j                  j
                        S r   )
_enum_attrr"   r^   ComputationModekComputeEigenvectorskNoEigenvectors)computes    r6   _eig_compute_attrr    sA    	9@fjj  55
 ::%%55
 r8   c                  |r|t         j                  k7  rt        d      | j                  \  }| j                  d   }|j
                  d d }|j                  t        j                  k(  xs |j                  t        j                  k(  }t        |j
                  d d |j                        }	t        |j
                  t        j                  |j                              }
t        |t        j                        }|	|
|
|g}|r|	g|}t        j                  d|j                        }t!        ||      } || |t#        |      t#        |            ^ }}}}|rt%        j&                  |d   |d         n|d   }t)        j*                  |t)        j,                  | dt        |t        j                  t        j                                    d	d
      }t/        | ||||      }|g}|r8| j                  t1        |         }t/        | ||||      }|j3                  |       |r8| j                  t1        |         }t/        | ||||      }|j3                  |       |S )Nz3Only the lapack implementation is supported on CPU.r   r   r   geev_ffir  )compute_leftcompute_rightr   rS  rT  )rH   rQ   r1  rU  rP  r   r   r   float32float64r   r   r  rW  r"   rV  rX  r  r%   complexr   rZ  r[  r\  r   r   )rK  r]  rS   rT   rU   r^  r_  r   realeigvals_avaleigvecs_avalra  rP  r`  rb  ro   vlvrrc  rd  outputr  s                         r6   _eig_cpu_loweringr    s   *;*B*BB
J
KK,,-,]]1(!!#2&*			rzz	)	M\-?-?2::-M$\//4l6H6HI,\//$55l6H6HIK,*bhh/)\<C)	*	*I**:|7I7IJ+	kY	?$#w'89R'S(9:T(UW1b"d "&ckk!A$!1Q4!
DQJ@R(ST
H" sJAx@!3&==V%D	!#z2r4	@B
MM"==V%D	!#z2r4	@B
MM"	-r8   c                P   | j                   dk(  r)t        j                  |t        j                  |            S | j                  dk(  | j                  t
        j                  k(  z  }t        j                  | j                  t              t        | j                        dz
        }|dz  dk(  }dgt        |j                        z  }d|d<   t        j                  |t        j                  |      |      }d|d<   t        j                  |t        j                  |      |      }t        j                  t        j                   t        |j                        t
        j"                  	      d
      }t        j$                  ||j                  |      }t        j$                  ||j                  |      }t        j&                  ||z  ||      }t        j&                  |||       }	t        j&                  |t        j                  |      |	      }	t        j                  ||	      S )Nr   r   )axisrX   r   r   r   )r   r   r   r   )r   r   r   rz  r   broadcast_dimensions)sizer   r  zeros_like_arrayimagr   nanr   cumsumastypeintr   r   pad_zerodeletearangerW  broadcast_in_dimrr  )
ro   r  is_realconj_pair_startpadsvr_shifted_leftvr_shifted_rightr  reims
             r6   _unpack_conjugate_pairsr    s   
 VVq[;;r3//344ffkaff./' !'''(9(9#(>-0\A-=?/#a'1,/
BHH	&$$r(GGB		"t4/$r(WWR25	299S]"((;R	@$  "((N'(("((>BD/
zz'O+R1AB"
zz/?RC8"
zz'3//3R8"	R	r8   c          	     	   | j                   \  }|j                  d d }|j                  dd  \  }}	||	k(  sJ |j                  }
t        j                  |
d      }|
t        j
                  t        j                  fv rd}n3|
t        j                  t        j                  fv rd}nt        d|
       |dk(  xr t        xr t        j                         dk\  }||r|r|t        j                  k(  r[|st        d      |rt        d	      | d
}t!        |||fz   |
      t!        ||fz   |      t!        |||fz   |
      t!        |||fz   |
      t!        |t        j"                        g}t%        ||      } || |||      \  }}}}}|r*t'        j(                  t*        d      }|rO| j-                  d t!        ||fz   |      t!        |||fz   |
      gt!        |||fz   |      g      } ||||      \  }|r| j-                  d t!        ||fz   |      t!        |||fz   |
      gt!        |||fz   |      g      } ||||      \  }njt.        j0                  j2                  }||t        j4                  k(  rdnd}t7        j8                          |r9| d}|
t        j
                  k(  rt        j                  nt        j                  }n/| d}|
t        j                  k(  s|
t        j                  k(  sJ |
}t!        ||fz   |
      t!        |||fz   |      t!        |||fz   |      t!        |t        j"                        g}|rt!        ||fz   |
      g|z   }t%        ||      } || ||||      ^ }}}}|r#t;        |      dk(  sJ t=        j>                  | }nt;        |      dk(  sJ |d   }t'        j@                  | dt!        |t        j"                              }t'        jB                  ||dd      }t!        ||fz   |      }tE        | ||||      }|g}|r1t!        |||fz   |      }tE        | ||||      }|jG                  |       |r1t!        |||fz   |      }tE        | ||||      }|jG                  |       |S )Nr                 ?TFzUnsupported dtype: r   i-  zANonsymmetric eigendecomposition requires cusolver 11.7.1 or newerz/Left eigenvectors are not supported by cusolversolver_geev_ffir  )leftrightr  	primitiverU  rP  onoffhybrid_eig_realhybrid_eig_comp)rJ   r  r  rX   r   r   rS  rT  )$rU  r   r   r   result_typer  r  	complex64
complex128r1  r   cusolver_get_versionrH   rO   RuntimeErrorNotImplementedErrorr   rW  rX  r   	lower_funr  r  r   gpu_use_magmar5   rP   r    initialize_hybrid_kernelsr   r%   r  r[  rZ  r\  r   )rK  r]  rS   rT   rU   r   r^  r   r  mr   complex_dtyper  have_cusolver_geevr`  rP  rb  r  ro   r  r  rc  unpackr  rJ   zerosrd  w_avalr  vl_avalvr_avals                                  r6   _eig_gpu_loweringr    s    ,,-,!!#2&*			BC	 $!Q	
a-


%..+-
rzz2::&&Gr}}--G
*5'2
33 D  8
8

,
,
.%
7  !3'*333
M  !
;= ='(8KJ!Q'/J!%}5J!Q'/J!Q'/J)I  yADc71J$>@Aq"b$~~5Nf	"++J!-}=J!Q/7 #:A#6FG  
 Wa$	#++J!-}=J!Q/7 #:A#6FG  
 Wa$  &&E!$(9(?(??dUe((*)*/:k&+rzz&9bllr}}m)*/:kbll"er}}&<<<m 	J!%u-J!Q'7J!Q'7J)	I zQD0%89IEiyADC!:"<>QB Vq[[
++q/aVq[[
A$a


c1k*bhh&G
H%eT84"zQD(-8&sJAv>!3&*1v-}=G	!#z2r7	CB
MM"*1v-}=G	!#z2r7	CB
MM"	-r8   c                   |s|rt        d      | \  }|\  }t        |d|      \  }}|gt        ||j                  |j                              t        |      z  j                  d      gfS )NzThe derivatives of non-symmetric eigenvectors are not supported. Only first-order derivatives of eigenvalues are supported. See https://github.com/jax-ml/jax/issues/2748 for discussion.F)rS   rU   r   )r  r^   _solver  r   _Tsum)	rA  rB  rS   rT   rU   rs   dar9  rn   s	            r6   eig_jvp_ruler    sw    "<
	DE E "!#"	Q%	O$!Q
q"))AGG,-15::2>?	??r8   r^   c               j    | d   | d   k7  rt        d|        | d   }||n
|d   |d   z
  }||f|ffS )Nr   r   zPArgument to symmetric eigendecomposition must have shape [..., n, n], got shape r   r0  )r   ri   r  r  r   s        r6   _eigh_shape_ruler    sh    
1Xr
	G	  Ah!#qqOA.. 
Q!r8   c                0    | t        j                  |       fS r   r   _complex_basetyper   r  s     r6   _eigh_dtype_ruler    s    	%%e,	,,r8   c                  ~| j                   \  }| j                  \  }}	|j                  d   }
||d|
fk(  st        d      |j                  d d }|t        j
                  k(  rt        d      |t        j                  k(  r|dk(  rt        d      |dk(  r|j                  }t        j                  |t        j                        rdnd	}t        j                  | d
|j                        }t        j                  t        d            t        j                  t        |rdnd            d}n9| d}|d}n|t        j                  k(  rdnd}|t        j                  |      d}t!        |t        j"                        }||	|g}t%        ||ddi      } || |fi |\  }}}t'        j(                  | d|      }t'        j*                  ||dd      }t-        | ||||      }t-        | ||||	      }||gS )Nr   r   z,subset_by_index not supported on CPU and GPUr   z,QDWH implementation is only supported on TPUr   z-Jacobi implementation is not supported on CPUhesyevd_ffiVrD  U)moderR  solver_syevd_ffirX   r   )rg   rl   rO  rS  rT  )rU  rP  r   r  r`   rf   re   r   r   r   r   r   r"   rV  uint8ordr   rW  rX  r   r[  rZ  r\  )rK  r]  rg   rh   ri   rl   r   r^  v_avalr  r  r   r   r   r`  r  algo_intra  rP  rb  rn   ro   rc  r  rd  s                            r6   _eigh_cpu_gpu_loweringr	    s    ,,-,==.&&!

!_A%>
L
MM!!#2&*$)))
L
MM$+++0Be0K
M
NN5 E&&ub.@.@ATtF,,xw-?-9-?-?AK hhs3x hhs%3S12F
 (((89Kh#5#<#<<!h288H+=>F*bhh/)vy))	kY67V
=$C+F+*!Q


c1i
0%eT84"sJAv>!sJAv>!
Q-r8   c          	     *   | \  }|j                   d   }||d|fk(  st        d      |\  }t        j                  t	        |      ||||      \  }	}
|
j                  |j                        }t        j                  |j                  ||f      }t        j                  d      5  t        j                  ||dt        j                  d d f   z   |dt        j                  f   z
  d      |z
  }d d d        t        |j                  dk(  rt        j                   nt        j"                  t        j$                  j&                        } | |t)        |	      |      |	      } ||	|z        }t+        |j,                        }|	|
f||ffS # 1 sw Y   xY w)	Nr   r   z8Derivatives not defined for partial eigen decomposition.rk   allow.rX   r<  )r   r  rm   r>   r;   r  r   r   r7  r   numpy_rank_promotioninteger_powr   newaxisr   r   dotr>  r?  r@  _H_extract_diagonalr  )rA  rB  rg   rh   ri   rl   rs   r  a_dotrn   w_realro   eye_nFmatr  vdag_adot_vdvdws                     r6   _eigh_jvp_ruler    sx    
$1ggbk!

!_A%>
B  &%kkm'%  )!V mmAGG!
((177QF
#%""7+ [??51S"**a%7#881S"**_;MMrRUZZD[ 	166Q;C,<,<--//	1#C1u%q)+
1d[ !"))*"
Vr2h	[ [s   $A	F		Frp   c                V    | d   | d   k7  rt        d|        | | d d | d   dz
  fz   fS )Nr   r   zHArgument to Hessenberg reduction must have shape [..., n, n], got shape r   r   r0  )r   r  s     r6   _hessenberg_shape_ruler    sQ    
1Xr
	G	  
cr
eBi!m--	--r8   c                
    | | fS r   rR   r  s     r6   _hessenberg_dtype_ruler  #      	r8   c                   | j                   \  }|j                  d d }|j                  d   }t        j                  |      st	        d|j                   d      t        j                  d|j                        }g | j                  t        |t        j                        }t        ||ddi      } || |t        j                  d      t        j                  |      	      \  }}}	t        j                  |	t        j                  | dt        |t        j                  t        j                                    d
d      }
t!        | ||
|| j                  d         t!        | ||
|| j                  d         gS )Nr   r   zKhessenberg requires the last dimension of a to be constant, got a.shape of r	  	gehrd_ffir   rO  r   )lowhighrS  rT  )rU  r   r   r   r1  r"   rV  r   rP  r   r   rW  rX  r   rZ  r[  r\  )rK  rs   a_avalr   r  r`  rP  rb  rw   rc  rd  s              r6   _hessenberg_cpu_loweringr$  '  sE   LL'&||CR *ll2!			a	 
 112	< = =**;E+AA{:rxx@A)	kY67V
=$sA288A;RXXa[A-!T4
DQJ@R(ST
H" sJAs}}Q7GHsJD#--:JK
 r8   rt   c                    | \  }}||k  rt        d|        |d   }|t        j                  ||      kD  rt        d      | S )Nz`The first argument to householder_product must have at least as many rows as columns, got shape r   z}The second argument to householder_product must not have more rows than the minimum of the first argument's rows and columns.)r1  r   min_dim)a_shape
taus_shaper  r  r  r   s         r6   _householder_product_shape_ruler)  D  sh    	$!QU
	&&-Y	01 1 m!a	
	EF F 
.r8   c                
   | j                   \  }t        |j                        s"t        j                  | |j                        g}nd }t        j
                  dt        j                  |      g||gd|      }|j                  gS )N(ProductOfElementaryHouseholderReflectorsr   result_typesoperandsr,   result_shapes)rP  r   r   r   eval_dynamic_shape_as_tensorcustom_callaval_to_ir_typer   )rK  rs   rw   aval_outr/  ops         r6   _householder_product_loweringr5  R  s{    mm)(	8>>	*))#x~~>@M M0((234y!#" ))r8   c                  | j                   \  }}|dk(  rN|j                  }t        j                  |t        j
                        rdnd}t        j                  | d|      }n| d}t        |ddi      }	 |	| ||      S )Nr   unorgqr_ffisolver_orgqr_ffir   r  )	rU  r   r   r   r   r   r"   rV  rX  )
rK  rs   rw   r   r#  r  r   r   r`  rb  s
             r6   %_householder_product_cpu_gpu_loweringr;  b  s    ll)&!5 LLE&&ub.@.@ATtF,,xw-?GK'((89K	k1a&	I$	c1d	r8   rx   c                &   | j                   \  fd}t        j                  t              fdt        j
                        }t        j                  d      }dk(  r
dk(  r||| fS t        j                  dt              |||| f      S )z-Unblocked LU decomposition, as a rolled loop.c           
     J   |\  }}}t        j                  d      }t        j                  d      }t        j                  |j                  t
        j                        r5|d d | f   }t        |j                        t        |j                        z   }nt        |d d | f         }t        j                  t        j                  || k\  |t        j                  |t
        j                               d|j                        }	|j                  |    j                  |	      }|j                  | |	gf   j                  ||	| gf         }|j                  |	| gf   j                  || |	gf         }|| | f   }
|j                  d d | f   j                  t        j                  || kD  |
dk7  z  |d d | f   |
z  |d d | f               }|d d | d f   || d f   z  }|t        j                  |d d d f   | kD  |d d d f   | kD  z  |t        j                   |            z
  }|||fS )NrW  r   )r  index_dtype)r   iotar   r   r   r   r   rq  r  r  argmaxrr  	full_likeinfr  r  _zeros)r   statepivotpermrs   m_idxn_idxt	magnituder  r?   a_outerr  r  s               r6   bodyz_lu_unblocked.<locals>.body  s   NE4HHWa EHHWa E""4"45
AqD'aaff+AFF+ia1g,i

3::eqj)S]]9rvvg5VWu{{	4AHHQKOOAE	aVW!QFG*%A77Aq67aVW.D 	
!Q$A	QT
szz519a"8!AqD'A+qAwOPA 1d
ma4j(G	CJJag*uT1W~/AB

7 35 	5A$>r8   r   rz  rW  )	r   r   fullminr   rW  r?  r   	fori_loop)rs   rL  rE  rF  r  r  s       @@r6   _lu_unblockedrP  |  s    	
$!Q4 ((C1I<"((
3%	'1	$!VQ 4			3q!9dUD!4D	EEr8   c           
        | j                   \  }}t        ||      }t        j                  |fdt        j
                        }t        j                  d|      }t        d||      D ]  }t        ||z
  |      }t        | |d|||z   f         \  }	}
}|j                  |||z    j                  |	|z         }|j                  |d j                  ||
|z            }| j                  |dddf   j                  | |
|z   ddf         } | j                  |d|||z   f   j                  |      } ||z   |k  s| j                  |||z   ||z   df   j                  t        | |||z   |||z   f   | |||z   ||z   df   ddd            } | j                  ||z   d||z   df   j                  t        j                  | ||z   d|||z   f   | |||z   ||z   df   t        j                  j                               }  | ||fS )z.Blocked LU decomposition, as an unrolled loop.r   rz  rW  NTr   rg   r   r<  )r   rN  r   rM  r   rW  r?  r   rP  r  r  r   addr  r?  r@  )rs   
block_sizer  r  r   rE  rF  r   r   block_pivot
block_permlu_blocks               r6   _lu_blockedrX    s"   	
$!Q	!Qi!
((A4"((
+%	'1	$Az" JaAE:A(5aAacE	l(C%KXHHQqsOa0E7712;??4
Q/0D	QRU*q.!+,-A	QR1Q3YH%A1uqy
$$q1uacd{


1QqS5!AaC%<!AacE1Q34K.D#4	9:a $$qstQqSTz


!A#$!A#+!AaC%1+25--2G2G!I  I JaJ 
E4r8   c                    | j                   dd }t        }t        t        |            D ]  }t	        j
                  |      }  ||       S )zCDefault LU decomposition in Python, where no better version exists.Nr   )r   rX  r   r   r
   vmap)r?   r   fnr  s       r6   
_lu_pythonr\    sF    wws|*"Z! a	"B 
A,r8   c                B    | \  }}| t        j                  ||      f|ffS r   r   r&  r   r  r  s      r6   _lu_shape_ruler`    s(    	$!Q	a#%t	++r8   c                    | t        j                  t        j                        t        j                  t        j                        fS r   r   r   r   rW  r  s     r6   _lu_dtype_rulerc    s)    	RXX&RXX(>	>>r8   c                H   t        j                  |      }t        |      dk(  sJ |\  }}t        j                  |      }t        ||      }dgdz  }d||z
  df|d<   t        j                  | d      }	t        j                  t        | d d d |f   d      |	|      }
|
t        j                  |||f      z   }
t        j                  t        j                  |||z
  ||z
  f      |	|ddf|ddff      }dgdz  }d||z
  df|d<   t        j                  t        | d |d d f         |	|      |z   }t        |
||   dddd      }t        ||ddd	      }t        j                  d
      5  |
t        |d      z  }t        |      |z  }d d d        ||z   S # 1 sw Y   z   S xY w)NrX   r  r   r   r   TF)r   r   rg   r   r;  highest)r   r   r   r   r   rN  r6  r  r<   r7  _triur   r   default_matmul_precision)r|   r  permutationr'  r  r  r   r   	l_paddingzeror9  u_eye	u_paddingr   lalaul_dotu_dots                     r6   _lu_jvp_innerrq    s    HHUO'	W			$!Q
))E
%	!Qi!kAo)a!eQ-)B-	B	$	ggeBq"1"uIr"D)4!#((51a&
!!!
''#((51q5!a%.14q!9q!Qi(*%kAo)a!eQ-)B-	ggeBrr1uIi058!5-$)TK"B%U$	&#&&y1 c2E#JNE 
 
s   )FF!c                P   | \  }|\  }t         j                  |      \  }}}t        }t        j                  |      d d D ]  }t        j                  |      }  ||||      }	|||f|	t        j                  j                  |      t        j                  j                  |      ffS Nr   )
rz   r>   rq  r   r   r
   rZ  r	   Zerofrom_primal_value)
rA  rB  rs   r  r|   r   rh  
lu_dot_funr  lu_dots
             r6   _lu_jvp_rulerx    s    "!&% IIaL"fk*88A;s &a*%J&b%-&
fk	"VW\\-K-KF-S%,\\%C%CK%P%R 
R Rr8   c                  | j                   \  }| j                  \  }}}|j                  d d }t        |t	        j
                  t        j                              }|j                  d   |dk(  r!t        j                  d|j
                        }	n| d}	t        |	|||gddi      }
 |
| |      \  }}}t        j                  |t        j                  | d|            }t        j                  |t        j                  | d|      dd	      }t        | ||||      }| j!                  d |g|g
      }t        j"                  fdd      } |||      \  }|||gS )Nr   r   	getrf_ffisolver_getrf_ffir   rO  r   GErT  r  c                    t        |       S r   )r   )r?   r  s    r6   r}  z&_lu_cpu_gpu_lowering.<locals>.<lambda>  s    %=a%C r8   Fr  )rU  rP  r   r   r   r   rW  r"   rV  rX  r%   subtractr   r[  rZ  r\  r  r  )rK  r]  r   r^  r_  
pivot_aval	perm_avalr   ra  r`  rb  r|   rE  rc  rd  r  perm_fnrF  r  s                     @r6   _lu_cpu_gpu_loweringr    s`   ,,-,$'MM!(J	!!#2&**bhhrxx&89)!5 ,,[,:L:LMK'((89K	k)1:y(I67V
=$ g&/"eT ,,ud11#q*E
F%d11#q)D
H"ZRB"KK$*#,+  /'NNC,13'
'5
!%$
eT	r8   c                   t        j                  | j                  d         t        j                  | j                  d         t        j                  | j                  d         g}t        d | j                  D              r8| j                  D cg c]"  }t        j                  | |j
                        $ }}nd }t        j                  d||g|      }|j                  S c c}w )Nr   r   rX   c              3  H   K   | ]  }t        |j                           y wr   r   r   )r   rs   s     r6   r   z(_lu_tpu_lowering_rule.<locals>.<genexpr>!  s     ?Aqww'	'?    "LuDecomposition)r-  r.  r/  r   r2  rP  r  r0  r   r1  results)rK  r]  r-  rs   r/  r4  s         r6   _lu_tpu_lowering_ruler    s    q)*q)*q)*,, 	??? }}
 ''QWW5M  MY	!"
 
s   'C"r|   tpuc                    t        | |||      S )zLU solve with broadcasting.)	_lu_solve)r|   rh  r   transs       r6   lu_solver  8  s     
2{Au	--r8   c           	     4   | j                   d   }t        j                  ||t        j                  |j                   dd        f      }|dk(  r)||d d f   }t        | |ddd      }t        | |dd      }n|dk(  s|dk(  rl|dk(  }t        | |ddd|      }t        | |dddd|	      }t        j                  |t        j                  d
|j                   d               \  }}||d d f   }nt        d|       t        j                  ||j                         S )Nr   r   TrR  F)r   rg   rX   )r   rg   r   r   )r   rg   r   r   r   rW  z&'trans' value must be 0, 1, or 2, got )	r   r   reshapemathprodr   sort_key_valr?  r1  )	r|   rh  r   r  r  r?   conjr  inds	            r6   _lu_solve_corer  >  s   hhqk!	kk!a17712;/01!
aZ	+q.AQ$d$OAQ$e<AzUaZA:DQ$e%)	+AQ$d$%)t	=Ak388G[=N=Nq=Q+RSFAs	#q&	A
=eWE
FF	Q	  r8   )   )static_argnumsc                p   t        | j                        dk  s| j                  d   | j                  d   k7  r$t        dj                  | j                              t        |j                        dk  r$t        dj                  |j                              | j                  |j                  dz   k(  }|rd|j                  d   | j                  d   k7  r/t        dj                  | j                  |j                              |dt
        j                  f   }nN|j                  d   | j                  d   k7  r/t        d	j                  | j                  |j                              t        j                  | j                  d d |j                  d d |j                  d d       }t        | g || j                  dd        } t        |g ||j                  d         }t        |g ||j                  dd        }t        }|D ]  }t        j                  |d
      }  || |||      }|r|d   S |S )NrX   r   r   zClast two dimensions of LU decomposition must be equal, got shape {}r   z*b matrix must have rank >= 1, got shape {}zWhen LU decomposition matrix and b have the same number of dimensions, last axis of LU decomposition matrix (shape {}) and b array (shape {}) must match.zWhen LU decomposition matrix and b different numbers of dimensions, last axis of LU decomposition matrix (shape {}) and second to last axis of b array (shape {}) must match)r   r   r   N)in_axesr   )r   r   r1  formatr   r   r  r   broadcast_shapes_broadcast_tor  r
   rZ  )	r|   rh  r   r  
rhs_vectorbatch_shaper[  r  r?   s	            r6   r  r  R  s   ]Q"((2,"((2,6
 $$*F288$46 6\A
AfQWWo' '
 ww!&&1*$*wwr{bhhrl" M rxx13 3 	
#rzz/Awwr{bhhrl" / rxx1	3 3 $$RXXcr]K4E4Ecr4JAGGTWUWLY+R7;7"#78"k+P[+P+:K:KB:O+PQ+A44qwwrs|45!" /a	"o	.B/[!U#! 6'a'r8   c                    ||    }||    }||   }|j                   |    j                  |      }|j                   |   j                  |      S r   )r  r  )r  rh  swapsjr?   rt  s         r6   _lu_pivots_body_fn_innerr  {  sP    Ah!!n!!n!q!%%a(+				q	!!r8   c                    |\  }}|j                   d d }t        }t        t        |            D ]  }t	        j
                  |dd      }  || ||      |fS )Nr   )Nr   r   r   )r  out_axes)r   r  r   r   r
   rZ  )r  permutation_and_swapsrh  r  r   r[  r  s          r6   _lu_pivots_body_fnr    s_    ,+u{{3B*"Z! 8a	"lQ	7B8	A{E	"E	))r8   c                    t        | j                        dk\  sJ | j                  dd }t        j                  |       j                  }|j
                  dd }|j
                  d   dk7  r1t        d|j
                  d    dt        j                  |              |j                  |dz         }| j                  d   }|}t        j                  t        j                  ||fz   t        |      |      }|d	k(  s|d	k(  r|S t        |      r$t        j                  |t        j                        n|}	t        j                  ||       \  }} t        j                   t        j                  d	t        j                        |	t"        || f      \  }
}|
S )
a  Converts the pivots (row swaps) returned by LU to a permutation.

  We build a permutation rather than applying `swaps` directly to the rows
  of a matrix because lax loops aren't differentiable.

  Args:
    swaps: an array of shape (..., k) of row swaps to perform
    permutation_size: the size of the output permutation. Should be >= k.
  Returns:
    An int32 array of shape (..., m).
  r   Nr   z3The last dim of swaps should be unsharded but got: z
 for type r   r  )out_shardingr   )r   r   r   typeofr  r  r1  r  r   broadcasted_iotar   rW  r   r~  rB   r   rO  r  )r  r~   r   swaps_shardingr  permutation_shardingr   r  rh  upperr   r  s               r6   !_generic_lu_pivots_to_permutationr    su    
U[[	Q		{{3B*;;u%...""3B'*$
	#$Jt{{5/A.B	DE E (..J4H.Ikk"o!!$$hh
aT!3z?')+ !VqAv#21#5"((1bhh
1%11+uE+u$$RXXa%:E%7+u9MO)&!	-r8   c               <    | \  }||k\  st        d| d|       |fS )NzOutput permutation size zE has to exceed the trailing dimension of the pivots. Got pivots size r0  )r   r~   pivots_sizes      r6   $_lu_pivots_to_permutation_shape_ruler    sD    ,+	[	(

"#3"4 5==HM	KL L 	r8   c               8    ~t        | ddd      } || |      S )N_lu_pivots_to_permutationr   F)num_non_batch_dimscolumn_majorrX  )rK  r   r~   r   rb  s        r6   &_lu_pivots_to_permutation_gpu_loweringr    s0     	!3 44MN12
H$	c6	r8   r   r   c                :    t         j                  |       \  }}||fS )aN  Computes the QR decomposition of a matrix.

  Args:
    a: an ``[..., m, n]`` batch of matrices, with floating-point or complex type.
  Returns:
    An ``(a, taus)`` pair where ``r`` is in the upper triangle of ``a``,
    ``q`` is represented in the lower triangle of ``a`` and in ``taus`` as
    elementary Householder reflectors.
  )geqrf_pr>   )rs   a_outrw   s      r6   geqrfr    s     Q+%	r8   c                >    | \  }}| t        j                  ||      ffS r   r^  r_  s      r6   _geqrf_shape_ruler    s$    	$!Q	a#%	%%r8   c                
    | | fS r   rR   rz  s    r6   _geqrf_dtype_ruler    r  r8   c                   t        j                  | j                  d         }t        j                  | j                  d         }||g}t        d | j                  D              r8| j                  D cg c]"  }t        j                  | |j
                        $ }}nd }t        j                  d||gd|      }|j                  S c c}w )Nr   r   c              3  H   K   | ]  }t        |j                           y wr   r  )r   r3  s     r6   r   z'_geqrf_lowering_rule.<locals>.<genexpr>  s$      	( x~~.	. 	(r  Qrr,  r  )rK  r]  ts_typer_typer-  r3  r/  r4  s           r6   _geqrf_lowering_ruler    s      q!12'a 01&6", 	(	( (  	))#x~~>M 
 M
y!" 
s   3'Cc                   | j                   \  }|dk(  r!t        j                  d|j                        }n| d}t	        |ddi      } || |      S )Nr   	geqrf_ffisolver_geqrf_ffir   r  )rU  r"   rV  r   rX  )rK  rs   r   r^  r`  rb  s         r6   _geqrf_cpu_gpu_loweringr    sW    ,,-,5 ,,[,:L:LMK'((89K	k1a&	I$	c1r8   r  rV   c               v    t        j                  | |      \  } }t        j                  | ||      \  }}}|||fS )a>  Computes the column-pivoted QR decomposition of a matrix.

  Args:
    a: a ``[..., m, n]`` batch of matrices, with floating-point or complex type.
    jpvt: a ``[..., n]`` batch of column-pivot index vectors with integer type,
    use_magma: Locally override the ``jax_use_magma`` flag. If ``True``, the
      `geqp3` is computed using MAGMA. If ``False``, the computation is done using
      LAPACK on to the host CPU. If ``None`` (default), the behavior is controlled
      by the ``jax_use_magma`` flag. This argument is only used on GPU.
  Returns:
    A ``(a, jpvt, taus)`` triple, where ``r`` is in the upper triangle of ``a``,
    ``q`` is represented in the lower triangle of ``a`` and in ``taus`` as
    elementary Householder reflectors, and ``jpvt`` is the column-pivot indices
    such that ``a[:, jpvt] = q @ r``.
  r  )r   rB   geqp3_pr>   )rs   jpvtrV   r  jpvt_outrw   s         r6   geqp3r    sB    " &&q$/'!T!,,q$),D%4	$	r8   c                @    | \  }}| |t        j                  ||      ffS r   r^  )r'  
jpvt_shaper  r  r  s        r6   _geqp3_shape_ruler  !  s&    	$!Q	*t||Aq13	33r8   c                    | || fS r   rR   )r   
jpvt_dtyper  __s       r6   _geqp3_dtype_ruler  %  s    	
E	!!r8   c               $   | j                   \  }}|dk(  r#t        j                  d|j                        }i }n?t	        j
                          t        j                  j                  }	| d}||rdnd}	d|	i}t        |ddd	
      }
 |
| ||fi |S )Nr   	geqp3_ffihybrid_geqp3r  r  rJ   r   r   r  r  )
rU  r"   rV  r   r    r  r   r  r5   rX  )rK  rs   r  rV   r   r#  r  r`  paramsrJ   rb  s              r6   _geqp3_cpu_gpu_loweringr  (  s    ll)&!5 ,,[&,,GKF((*  &&E'(5KdUeuF	kaA,	O$	c1d	%f	%%r8   r  )r   r  c               h    | \  }}|r|nt        j                  ||      }|r
||f||f|ffS ||f||ffS r   r^  )r   r   r   r  r  r  r   s          r6   _qr_shape_ruler  =  sJ    	$!Qadll1a0!#+1a&1a&1$	A1a&1a&1AAr8   c               Z    |r&| | t        j                  t        j                        fS | | fS r   rb  )r   r   r  s      r6   _qr_dtype_ruler  B  s'    3;%RXX.	/O%Or8   c                  | \  }|\  }t         j                  ||d|      ^}}}	|j                  ^ }
}}||k  s|r||k7  rt        d      |r
|d|	d   f   }t	        ||      }t        |      |z  }t        |d      }|t        |      z
  }t        j                  t        j                  |j                  ||f      t        |j                  dz
              }||||j                  j                  |j                        z
  z  z   }|||z
  z  |z   }||z
  |z  }|r/t        j                   j#                  |	d         }|||	d   f|||ffS ||f||ffS )NFr   z1Unimplemented case of QR decomposition derivative.r   r   rX   )r   r>   r   r  r   r  r<   r   r   r7  r   r   r   r  r  r	   rt  ru  )rA  rB  r   r   rV   r?   dxr   r   r   r  r  r  dx_rinv
qt_dx_rinvqt_dx_rinv_lowerdoIdqdrdps                        r6   qr_jvp_ruler  E  sh   "!#"YYq85IYV(!QWW(1aU}a
9; ;	C1IBQ#'!uw*:r*"-.."	oochhrxx!Q0%
!8K2LM!	Ajoo44Z5E5EFFGG"BOw&"R1"		'	'!	-Bq!A$<"b"%%
Q"b	r8   c                  | j                   ^ }}}|dk(  s|dk(  r|r|nt        j                  ||      }t        j                  t        j
                  | j                  ||f      g |||t        |      t        |      dz   f      }t        j                  g |||d| j                        }	|rCt        j                  g ||dt        j                  t        j                              }
||	|
fS ||	fS |rVt        j                  g ||dt        j                  t        j                              }t        | ||      \  }	}
}|
dz  }
nt        |       \  }	}||k  rt        |	dd |d |f   |      }nm|rSdgt        |      dz   z  d||z
  dfgz   }t        j                  |	t        j                  |	      |      }t        ||      }nt        |	|      }|	dd |d |f   }	t!        |	      }	|r||	
fS ||	fS )Nr   r   rz  r  .r  )r   r   r&  r   r  r7  r   r   rM  r   rW  r  r  rx   r  r  rf  )rs   r   r   rV   r   r  r  r   r   r   r   r  rw   r  s                 r6   _qr_loweringr  ^  s   gg:q!!VqAvQ 2ASXXagg1v60z010a0!*os:/BC	EA 	$:$q$!$aqww7A
((#Z##Qbhhrxx.@
Aa1Wna4K88$j$!$arxx/ABDq$)4JAq$FAAhGAtUAc2A2rrkND1A;#j/A-.1a!eQ-@D399Q<&AAt$AAt$A	#rr2A2+AAh!a7N	
A+r8   ra   c               H    | d   | d   k7  rt        d|  d      |r| | fS | fS )Nr   r   z6The input to schur must be a square matrix. Got shape r	  r0  )r   r   r  s      r6   _schur_shape_ruler    s@    
1Xq

@qIK K0%>uh>r8   c                   |r| | fS | fS r   rR   )r   r   r  s      r6   _schur_dtype_ruler    s    0%>uh>r8   c               b   ~|rt        d      | j                  \  }|j                  d d }|j                  t        j
                  k(  xs |j                  t        j                  k(  }t        j                  d|j                        }t        |t	        j                  t        j                              }	t        |j                  d d |j                        }
|r	|||
|
|	|	g}n|||
|	|	g}|r$t        j                  j                  j                  n#t        j                  j                  j                  }t        ||ddi      } || |t!        |      t!        t        j                  j"                  j$                              ^}}}}t'        j(                  |t'        j*                  | dt        |t	        j                  t        j                                    dd	      }t-        | |||| j.                  d         }|g}|r-t-        | |||| j.                  d
         }|j1                  |       |S )Nz=The sort feature of LAPACK's gees routine is not implemented.r   gees_ffir   r   rO  )r  sortrS  rT  r   )r  rU  r   r   r   r  r  r"   rV  r   rW  r   r  kComputeSchurVectorskNoComputeSchurVectorsrX  r  SortkNoSortEigenvaluesr   rZ  r[  r\  rP  r   )rK  r]  r   r   r   r^  r   r  r`  ra  r  rP  r  rb  
schur_formschur_vectorsr  rc  rd  r  s                       r6   _schur_cpu_loweringr    s   
GI I ,,-,!!#2&*			rzz	)	M\-?-?2::-M$**:|7I7IJ+*bhhrxx&89)\//4l6H6HI,	|\<I'I |\9iPI 
 ll""77<<''>> 
 
kY67V
=$(,	7D)fll''::;)=%*ma 
DQJ@R(ST
H" (ZZ(+a(8:*<&,S*b--0]]1-=?M
MM- 	-r8   r   c                   | \  }}t        j                  ||      }|3|r|d|fk7  rt        d      t        j                  ||d   |d   z
        }|r|f||r|n|f|r||ffS ||ffS |ffS )Nr   z4full_matrices and subset_by_index cannot both be setr   )r   r&  r1  )r   r   r   ri   r  r  r  r  s           r6   _svd_shape_ruler    s    	$!Q	a	$ QI5MNN<<oa0?13EEFD		
AD)q)  "&q)  78Or8   c               @    t        j                  |       }|r|| | fS |fS r   r  )r   r   r  
real_dtypes       r6   _svd_dtype_ruler    s*    $$U+*ue##;r8   r  )rl   c                  | \  }|\  }t         j                  |dd||      \  }}	}
|r=|r;t        j                  |j                  d   |j                  d         st        d      t        |	      t        |
      }}|dd d d f   }||z  |z  }t        |j                        }|s|f|ffS |t        |      z   |t        |      z
  z  }t        j                  |j                  |j                  d   |j                  d   f      }t        j                  |t        |j                  dz
              }d	||z   z  |z
  }|j!                  |j                        |z  }t        |j!                  |j                              |z  }|d
k(  j!                  |j                        }d	||z   z  |z
  }t#        |      }d|t        |      z
  z  |j!                  |j                        z  }|	|j!                  |j                        |t        |      z   z  |z   z  }||j!                  |j                        |t        |      z   z  z  }|j                  dd  \  }}||kD  r/||z  }|||	||z  z  z
  |j!                  |j                        z  z   }||kD  r8t        |      |	z  }||||
|z  z  z
  |j!                  |j                        z  z   }||	|
f||t        |      ffS )NFTr   r   r   zBSingular value decomposition JVP not implemented for full matrices.rX   r   r   g      ?)r   r>   r   definitely_equalr   r  r  r  r  r  r   r7  r   r   r   r   r  _construct_diagonal)rA  rB  r   r   ri   rl   AdAr   r  VtUtr  s_dimdSdss_diffss_diffs_zerosFdSSSdSs_zeross_inv	s_inv_mat	dUdV_diagdUdVr  r  dAVdAHUs                                  r6   _svd_jvp_ruler    s   
 "!#"ZZu  (!Q 
##AGGBK= JL L Q%Ba"
CqL/%	Bw{"!"	4";RY52e9#45'((177QWWR[!''"+$>?-//-w||a7G1HI-7]"#m3!QWW"#
5<< !B&#!VOOAGG$'
q7{
g
%%!%()BBK 9#3#3AGG#<<)AHHQWWr#w/);<"AHHQWWr#w/0"	
$!QU
q&C	sQ"s(^#u||AGG'<<	<BUb6A:D	ta29o%agg)>>	>B
Qb"bf%	%%r8   c               :   | j                   d d }| j                   dd  \  }}t        j                  |dz   dt        j                  | j                              }|s|fS |rbt        ||      }t        j                  t        j                  | j                  ||f      g |||t        |      t        |      dz   f      }n't        j                  |||fz   d| j                        }t        j                  |dz   d| j                        }	||k  r|	|}	}|||	fS )Nr   r   r   rz  r   )r   r   )	r   r   rM  r  r   maxr  r7  r   )
rs   r   r   r  r  r  r   r  r   rn   s
             r6   
_empty_svdr    s   +	
$!Q	hh{T!1C,A,A!'',JK!	4Kq!9DSXXaggd|<7{7D7$7!+.K0@10DE	GA 	1v%q8A	hh{V#Qagg6!UaqA	
Aq.r8   c                J    d}|d}| sd}t        |      S |sd}t        |      S )Nr  TNS
_char_attr)r   r   r  s      r6   _svd_computation_attrr$  %  s<    	$M	D 
D	 D	D	r8   c          
     b   | j                   \  }| j                  d   }|j                  dd  \  }	}
|j                  d d }||dt        |	|
      fk(  st	        d      |	dk(  s|
dk(  r% t        j                  t        d      | |||      S |dk(  rb||t        j                  k(  r!t        j                  d|j                        }n?|t        j                  k(  r!t        j                  d	|j                        }nt	        d
      t        ||      }t        |t!        j                  t         j"                              }|r| j                  \  }}}nx| j                  \  }t        g ||	|r|	nt%        j&                  |	|
      |j                        }t        g ||r|
nt%        j&                  |	|
      |
|j                        }|||||g}t)        ||ddi      } || ||      \  }}}}}nt+        | |||||      \  }}}}t        j,                  | dt        |t!        j                  t         j"                                    }t        j.                  ||dd      }t1        | ||||      }|g}|r7| j                  dd  \  }}t1        | ||||      }t1        | ||||      }|||gz  }|S )Nr   r   z/subset_by_index not implemented for CPU and GPUTr  r   r   r   	gesdd_ffi	gesvd_ffiz?The SVD Jacobi and Polar algorithms are not implemented on CPU.rO  )r  )r   r   r   rl   rS  rT  r   )rU  rP  r   rN  r  r   r  r  r   r   r"   rV  r   rd   r$  r   r   rW  r   r&  rX  _svd_gpu_sub_loweringr[  rZ  r\  )rK  r]  r   r   ri   r   rl   r^  s_avalr  r  r   r`  r  ra  u_avalvt_avalrP  rb  r  r   r   vtrc  r  rd  r   s                              r6   _svd_cpu_gpu_loweringr.  /  s    ,,-,==&			BC	 $!Q!!#2&*

!_C1I%F
O
PP!VqAv<4>>*t<#	  5 I)=)==..{L<N<NOk	loo	%..{L<N<NOk
KM M ];DJ(:;I #ffggf  HZ H H!.ADLLA4FH'--/f  Lj L"/QT\\!Q5GLIJL(..0g vvw	BIy891v?DS'5Aq!R*39F6@>P5>	@NAq"d 

c1k*bhhrxx>P&Q
R%eT84"sJAv>!3&mmAB'OFG j"a@A	!#z2r7	CB
q"gF	-r8   c                  | j                   \  }|r| j                  \  }}}	n&| j                  \  }t        d|j                        x}}	|j                  d d }
t        |
t        j                  t
        j                              }t        |
      }|j                  dd  \  }}t        j                  ||      }d}i }d}d}||t        j                  k(  r	 |dk(  xr |dk  xr |dk  }n+|t        j                  k(  rd}n|t        j                  k(  rd}d}|r| d}	 | xr |dkD  xr |dkD  }n |r	| d	}| }n| d
}| }||k  }d|i}|rd}|s|rJt        g |
||r|n||j                        }t        g |
||r|n||	j                        }|||||g}n|r|||	||g}n||||	|g}t        ||ddi|      } || |f| |d|\  }}}}}|s|r|rt!        j"                  |t%        j&                  t        j(                  t+        t-        |            |dz   |fz                     }t        j.                  |j                  t
        j0                        rOt!        j2                  t!        j4                  |      t!        j6                  t!        j8                  |                  }|s|st        |j                        }t%        j:                  | || j                  d   t        j<                  |gt
        j>                        |
||fz   t        j@                  |gt
        j>                              }t%        j:                  | || j                  d   t        j<                  |gt
        j>                        |
||fz   t        j@                  |gt
        j>                              }|r||||fS ||||fS # t        j                  $ r d}Y w xY w# t        j                  $ r d}Y w xY w)NrR   r   Fr   i   Tsolver_gesvdj_ffi    solver_gesvdp_ffisolver_gesvd_ffi
transposedr   )rP  rQ  r  r&  r   )start_indiceslimit_indicesstridesrX   )!rU  rP  r   r   r   r   rW  r   r   r&  r   r   InconclusiveDimensionOperationre   r   rX  r%   r   r   dense_int_arrayr~  r   r   r   r   r  r  negater  slice_opr  int64ones) rK  r]  r   r   r   rl   r^  r*  r+  r,  r   ra  nbr  r  r   r4  r  
use_jacobi	use_polarr  r`  econr  rP  rb  r  r   r   r-  rc  nds                                    r6   r)  r)  u  s   ,,-,!mmFFGmmGF"2|'9'9::FW!!#2&**bhhrxx&89)
:"			BC	 $!Q
ll1a!*&" *))|';';;%-I!t)IT	j L'''JL&&&I,'((9:K41r64a"fd '((9:KD'((89KD QJJ'Fl9 ;:;q;t!;V\\JF;:;q;t!;W]]KFvvvyAIvw	BIvvw	BI	kY67V+7
9$ 3 <4x&0<4:<!Q2tI:	
RXXeE"I&6"q&"&EFG
IB 
}}\''););<;;sxx|SZZ%=>b|!!"b
--Qa 0&(hhtRXX&>&0Aq6&9 "bhh 79a ==b#--"2')xxbhh'?'1QF':!#"rxx!8:b b!T>aT>C .. j .. ds$   N< O <OOO10O1r   c                V    | d   |d   k7  s|d   |d   k7  rt        d|  d| d      |S )Nr   r   znsymmetric_update expects a rectangular matrix of shape (m, n) and a square matrix of shape (n, n). Got shapes rj  r	  r0  )r'  c_shaper  s      r6   _symmetric_product_shape_rulerE    sO    QZ71:wqz!9
	55<IU7)1	NO O 
.r8   c                  t        j                  | g t        | j                  dz
        | j                  dz
  | j                  dz
        }|t        j                  | |t         j
                  j                        z  ||z  z   S )NrX   r   r<  )r   r   r   r   r>  r?  r@  )rs   r  r   r   a_Ts        r6   _symmetric_product_jax_fnrH    sv    aE5!,EaffqjE!&&1*EF#	!!--/ 
/15
: :r8   c                6   |j                   d d \  }}|j                  }t        d|      x}	}
t        j                  |||	      }t        j                  |||
      }t        j                  |  dddi      }|j                  |||	|
g      } ||||||d	      S )
NrX   rR   solver_syrk_ffir   r   r  )rU  F)r   )rU  r   r   r   r[  r   r  r  )r+   rK  a_tensorc_tensorr   r   r#  c_avalr   
alpha_aval	beta_avalalpha_array
beta_arrayrb  s                 r6   _symmetric_product_gpu_loweringrR    s    <<#.&&
,,%&r511*y##C
;+""3i8*			XJo623Q
9$ffj)DE#	c8X{J%	PPr8   )rX   rX   r   r   c                   | d   | d   k7  rt        d|  d      |rdnd}| d   ||   k7  rt        d|  d| d      |S )	Nr   r   zGThe first input to triangular_solve must be a square matrix. Got shape r	  r   r   z7Incompatible shapes for arguments to triangular_solve: rj  r0  )r'  b_shaper   r  
common_dims        r6   _triangular_solve_shape_rulerW  	  su    QZ71:
			  rB*R[GJ''

A'%)1	  
.r8   c                    | S r   rR   )r   r  r  s      r6   _triangular_solve_dtype_rulerY  	  s    	,r8   c               4   |j                   dd  \  }	}
rdnd}rt        | |       nt        | |      } t        j                  |       } rt        |       n| } r| j                         n| } t        | j                  dk(  rt        j                  nt        j                  t        j                  j                        }fd}rd| j                   dd  j                   dd  cxk(  r|	|	fk(  rn J |j                   dd  |	|
fk(  sJ |	|
kD  r | || |            S  | ||       |      S | j                   dd  j                   dd  cxk(  r|
|
fk(  rn J |j                   dd  |	|
fk(  sJ |	|
k  r | |||             S  || ||             S )Nr   r   r   r   rX   r<  c           	     (    t        |       S Nr   r   )rhsrs   r   r   rg   r   r   s    r6   	a_inversez/_triangular_solve_jvp_rule_a.<locals>.a_inverse)	  s     Asiu(3*79 9r8   )r   r<   rf  r   negr  r  r   r   r  r>  r?  r@  )g_aansrs   r   r   rg   r   r   r   r  r  r   r  r_  s     ` `````     r6   _triangular_solve_jvp_rule_arc  	  s~    
$!Qaa!!caRuSA##3C#!
s#388q=c.>.>--//	1#9 9 99RS>QWWRS\3aV3PP		"#1a&8PPP1us3}%%3%%99RS>QWWRS\3aV3PP		"#1a&8PPP1us3}%%in%%r8   c          	        t        j                  |      st        j                  |      sJ t        |       t        j                  u r#t        j                  |j
                        }d |gS t        || ||| ||      }d |gS r\  )r   is_undefined_primaltyper	   rt  r  r   )		cotangentrs   r   r   rg   r   r   r   cotangent_bs	            r6    _triangular_solve_transpose_ruleri  >	  s    
 ##A&2+A+A!+DD	D	)_$,,qvv&K 		 #1i9).O/:1>@K 	r8   c          	        | \  }}|\  }	}
|	t         j                  u r
|rft        j                  ||
d      }|j                  |j                  d d |j                  d   |j                  d   z  fz         }|j
                  dz
  }nst        j                  ||
d      }|j                  |j                  d d |j                  d   |j                  d   z  |j                  d   fz         }|j
                  dz
  }t        |||||||      }|j                  |j                        |fS t        d t        | |      D              }t        j                  ||	|      }t        j                  ||
|      }t        |||||||      dfS )	Nr   r   r   rX   r   c              3  D   K   | ]  \  }}||j                   |     y wr   r2  r   rI  r  s      r6   r   z2_triangular_solve_batching_rule.<locals>.<genexpr>a	  (      "tq!= 
 "    r   )
r   
not_mappedmoveaxisr  r   r   r   nextr   bdim_at_front)batched_argsr   r   rg   r   r   r   r?   rt  bxbyy_flatbdim_outout_flatr  s                  r6   _triangular_solve_batching_rulerz  M	  s~    
$!Q&"b8


Ar2
&ayy"qwwr{)B(DDEf!h


Ar2
&ayy"!''"+*CQWWR[)QQRf!h	6Ye[#%H AGG$h.. "s<'D " "Dq"d+Aq"d+AAqIU(3*79:;< <r8   c          
        | j                   \  }|r|st        j                  |      }d}|sd}	n|rdnd}	t        j                  ||t
        j                  j                  |      t
        j                  j                  |      t
        j                  j                  |      t        j                  j                  |	            }
t        j                  | |
|      gS )NFNO_TRANSPOSEADJOINT	TRANSPOSE)rP  r$   ConjOpr%   r   r#   rI  rJ  TransposeAttrr   lower_with_sharding_in_types)rK  rs   r   r   rg   r   r   r   r_  r   r   s              r6   _triangular_solve_loweringr  i	  s    mm)(AAK	I(	kIQ2;;??9#=[[__U3[[__]; ..229=	?# 
+
+Ch
?	@@r8   c          
        | j                   \  }}	|r|st        j                  |      }d}t        j                  |j                        t
        v rqt        j                  d|j                        }
d\  }}}t        |
||	g|ddi|      } || ||g|t        |      t        |      t        ||      t        |      dS |r|rdnd	}nd
}t        j                  ||t        j                   j#                  |      t        j                   j#                  |      t        j                   j#                  |      t        j$                  j#                  |            gS )NFtrsm_ffi)rR   rR   Tr   r   )rQ  batch_partitionable)siderR  trans_xdiagr}  r~  r|  )rU  r$   r  r   r   _cpu_lapack_typesr"   rV  rX  _matrix_side_attrrY  _matrix_transpose_attr_matrix_diagonal_attrr%   r   r#   rI  rJ  r  )rK  rs   r   r   rg   r   r   r   r#  b_avalr`  r   rN  r  rb  r   s                   r6   _triangular_solve_cpu_lowerr  }	  s=    <<.&&		!AKXXfll00,,ZFK-9*E:*!' =* =891v4GID Q ;E ;&y1&u-.{KH*=9	; ; *)i i  Ar{{y'A!#!7!#!?!$!2!2!6!6y!AC D Dr8   r   c                    t        || fi |S r   r]  )g_br  rs   r   kwss        r6   r}  r}  	  s    '73'F#'F r8   c                n    | d   | d   k7  s| d   dk(  rt        d|  d      | \  }}| |f|dz
  f|dz
  ffS )Nr   r   z<The input to tridiagonal must be a square matrix. Got shape r	  r0  )r   r  r  s      r6   _tridiagonal_shape_ruler  	  s]    
1XqU1X]

FugQOQ Q	$!Q	q1uhQ	((r8   c                8    t        j                  |       }| ||| fS r   r  )r   r  r  s      r6   _tridiagonal_dtype_ruler  	  s!    $$U+*	
J	--r8   c                  | j                   \  }| j                  \  }}}}|j                  d d }	|dk(  rs|j                  t        j
                  k(  xs |j                  t        j                  k(  }
|
rdnd}t        j                  | d|j                        }dt        |      i}n	| d}d|i}t        |	t        j                        }t        |g | j                  |d	d	i
      } || |fi |\  }}}}}t        j                  | d	|      }t        j                  ||dd      }t!        | |	|||      }t!        | |	|||      }t!        | |	|||      }t!        | |	|||      }||||fS )Nr   r   r  r  trd_ffirR  solver_sytrd_ffirg   r   rO  rS  rT  )rU  rP  r   r   r   r  r  r"   rV  rY  r   rW  rX  r   r[  rZ  r\  )rK  rs   rg   r   r#  arr_avald_avale_aval	taus_avalr   r  r   r`  r  ra  rb  arrr   erw   rc  r  rd  s                          r6   _tridiagonal_cpu_gpu_loweringr  	  s|   LL'&(+%(FFI||CR *5 <<2::%C)CDTtF,,xw-?NK'./F'((89KuF*bhh/)	8s}}8i8V
%$ sA00#q!T4


c1i
0%eT84" j"c8D#sJAv>!sJAv>!	!#z2tY	G$	aDr8   r   c                V    | |k7  s| |k7  rt        d      | |d d k7  rt        d      |S )NzKtridiagonal_solve requires that all diagonal arguments have the same shape.r   zrtridiagonal_solve requires that the leading ndim-1 dimensions of b equal the dimensions of the diagonal arguments.)r   )dl_shaped_shapedu_shaperU  r  s        r6   _tridiagonal_solve_shape_ruler  	  sL    H0
	  "
	:; ; 
.r8   c               B    | d}t        |ddi      } || ||||      S )Nsparse_gtsv2_ffir  r   r  r  )rK  r   r   r   r   r   r`  rb  s           r6   _tridiagonal_solve_gpu_loweringr  	  s3    %&&67+	k1a&	I$	c2q"a	  r8   c           	        ~| j                   d   }|j                  d d }t        j                  d|j                        }t        |t        j                        }	t        |g | j                   |	ddddd	      }
 |
| ||||      ^ }}}t        j                  | d|	      }t        j                  ||d
d      }t        | ||||      gS )Nr   r   gtsv_ffir   r   rX   r  )r   r   rX   r  rO  rS  rT  )rU  r   r"   rV  r   r   r   rW  rX  r   r[  rZ  r\  )rK  r   r   r   r   r  r  r   r`  ra  rb  r  b_outrc  r  rd  s                   r6   _tridiagonal_solve_cpu_loweringr  	  s    <<&||CR ***:v||D+*bhh/)	k(B#,,(B	(B9:qQ15M
O$ b!R+/1eT


c1i
0%eT84"
"3
Bv
F	GGr8   c                :   t        j                  ||j                  dz         |z  }|j                  ddd d d f   j	                  | ddd d f   |dd dd d f   z        }|j                  dd dd d f   j	                  |dd dd f   |ddd d d f   z        }|S )Nr  .r   r   )r   r  r   r  rS  )r   r   r   r   rt  s        r6   _tridiagonal_productr  	  s    	kk!QWWt^$q(!dd3A:2c12tm,qcrc1~=>!dd3Q;BsCRC~.3A:>?!	
(r8   c                   | ^ }}|^ }}t        j                  |  }t        d |D              r|}n<t        g t	        t
        j                  |      | }t        j                  ||       }t        j                  g || }	||	fS )Nc              3  R   K   | ]  }t        |      t        j                  u  ! y wr   )rf  r	   rt  )r   r   s     r6   r   z._tridiagonal_solve_jvp_rule.<locals>.<genexpr> 
  s     4QaGLL	 4s   %')r   r>   r   r  mapr   instantiate_zerosadd_tangents)
rA  rB  diagsr  	diags_dotb_dotrb  r^  
matvec_dotans_dots
             r6   _tridiagonal_solve_jvp_ruler  	  s    )5!9e  '*#4)44
C%Qs2+?+?'KQSQJ
//%*
-C$$1e1S1'	gr8   c           	     d   t        j                  |      s?t        j                  |      s*t        j                  |      st        j                  |      sJ t        |       t        j                  u r t        j                  |j
                        }nt        j                  t        j                  |ddd f         |dd df   f|j                  dz
        }t        j                  |ddd f   t        j                  |dd df         f|j                  dz
        }t        ||||       }d d d |gS )N.r   r   )r   re  rf  r	   rt  r  r   concatenater  r   r   )rg  r   r   r   r   rh  dl_transdu_transs           r6   !_tridiagonal_solve_transpose_ruler  
  s	    $$R(B,B,B1,E$$R(b.D.DQ.GH 
H	)_$,,qvv&K 4 4RRS\ BBsCRCxLQ!wwqy*H37S-A-A"S"1"W+-NO!wwqy*H#Ha9EK
dK	((r8   c                   | \  }}}}|\  }}}}	|t         j                  u r|t         j                  u r|t         j                  u rt        j                  ||	d      }|j                  |j                  d d |j                  d   |j                  d   |j                  d   z  fz         }
|j
                  dz
  }t        ||||
      }|j                  |j                        |fS t        d t        | |      D              }t        j                  |||      }t        j                  |||      }t        j                  |||      }t        j                  ||	|      }t        ||||      dfS )Nr   rk  r   rX   c              3  D   K   | ]  \  }}||j                   |     y wr   r2  rm  s      r6   r   z3_tridiagonal_solve_batching_rule.<locals>.<genexpr>$
  rn  ro  r   )
r   rp  rq  r  r   r   r   rr  r   rs  )rt  r   r   r   r   r   bdlbdbdubbb_flatrx  ry  r  s                 r6    _tridiagonal_solve_batching_ruler  
  sT   ,"aQ#r3	X   H	X   !R$AYYqwws|QWWR[1772;5N'OOPFvvzH QF3HAGG$h.. "s<'D " "D			C	.Bq"d+A			C	.Bq"d+ARB*A--r8   c                   d }t        j                  ||d   |d   z  |d   |d   z  f| dd  |dd  |dd  |dd d d f   fd      \  \  }}\  }}d }	t        j                  |	|||fdd      \  }
}t        j                  |
d    |fd      S )	Nc                b    | \  }}|\  }}}}||||z  z
  z  }|||z  z
  |||z  z
  z  }	||	f||ffS r   rR   )
carryargscpr  rs   r   r  r   cp_nextdp_nexts
             r6   fwdz(_tridiagonal_solve_jax_impl.<locals>.fwd-
  sX    FBJAq!Q1q2v:G1r6za!b&j)GWBx''r8   r   r   r1  )unrollc                $    |\  }}||| z  z
  }|| fS r   rR   )xnr  r  r  r?   s        r6   bwdz(_tridiagonal_solve_jax_impl.<locals>.bwd8
  s!    FB
R"WAb5Lr8   T)r  reverse)r   scanr   r  )r   r   r   r   r  r  finalr  r  r  endrb  s               r6   _tridiagonal_solve_jax_implr  ,
  s    ( &**	BqEAaDL!A$1+&AB1212!"a%(I*1ehr2
 sEB8BM(#s	#d)S)1	--r8   c                    t         }t        | j                  dz
        D ]  }t        j                  |      }  || |||      S r  )r  r   r   r
   rZ  )r   r   r   r   r  impls         r6   _tridiagonal_solve_jaxr  @
  sA    	$$1 a88D>D	b!R	r8   )r   r   r   rX   r   r   r   c                &   | j                   dk\  sJ |j                   dk\  sJ t        j                  | j                  d d |j                  d d       }t	        |      }t        | g || j                  dd        } t        |g ||j                  d         }| j                   dz
  g|j                   dz
  gft        t        |            t        t        |            ff}t        j                  | ||t        j                  j                        S )NrX   r   r   r   )dimension_numbersr=  )r   r   r  r   r   r  listr   dot_generalr?  r@  )rs   r   r  n_batchr  s        r6   _broadcasted_matvecr  ^
  s    	
1	
1$$QWWSb\1773B<@+'A44qwwrs|45!A22aggbk23!!qvvzl3d5>6JDQVW^Q_L`5ab	A1BcmmNcNc	ddr8   c                <   | j                   dk\  rh|j                   | j                   | j                   dz
  fv rA| j                  d   | j                  d   cxk(  r|j                  | j                   dz
     k(  s'n t        d| j                   d|j                         y )NrX   r   r   r   z]The arguments to solve must have shapes a=[..., m, m] and b=[..., m, k] or b=[..., m]; got a=z and b=)r   r   r1  )rs   r   s     r6   _check_solve_shapesr  j
  s    
&&A+!&&QVVQVVaZ$88
''"+
;
(;
;
	../ggYgaggY	HI I <r8   c           	     R    t         |       t        d t         j                  d d dz   |j                        D              }t	        j
                  ||t        |j                              }t        t	        j                               \  }t        t        j                   fdfdfd      } j                  |j                  dz   k(  r ||      S  t        j                  ||j                  dz
  t         j                  |j                        dz
        |      S )	Nc              3  4   K   | ]  \  }}|d k(  r|n|  yw)r   NrR   )r   d_ad_bs      r6   r   z_solve.<locals>.<genexpr>v
  s)      G S !8C, Gs   r   r  c                    t        |       S r   )r  )r?   rs   s    r6   r}  z_solve.<locals>.<lambda>
  s    #Aq) r8   c                "    t        |d      S )Nr   r  r  r  r?   lu_rh  s     r6   r}  z_solve.<locals>.<lambda>
  s    #{AQ? r8   c                "    t        |d      S )Nr   r  r  r  s     r6   r}  z_solve.<locals>.<lambda>
  s    8Caq#I r8   )solvetranspose_solver   )r  r   r   r   r   r  r   r   r|   stop_gradientr   r   custom_linear_solver
   rZ  r  )rs   r   	out_shaper  custom_solver  rh  s   `    @@r6   r  r  q
  s    a  G$'t(;QWW$EG G)	1iqvv7! 3,,Q/0#q+&&)?I	K,
 VVqvvz? G388L!&&1*c!&&!&&.AA.EFqIIr8   c                    t        j                  | g t        | j                  dz
        | j                  dz
  | j                  dz
        S )NrX   r   )r   r   r   r   r{   s    r6   r  r  
  s@    	qFE!&&1*-FqvvzF166A:F	GGr8   c                4    t        |       j                         S r   )r  r  r{   s    r6   r  r  
  s    	Ar8   c                $    | t        |       z   dz  S )NrX   )r  r{   s    r6   r;   r;   
  s    1r!u9/#9r8   c                    | j                   ^ }}}t        j                  t        ||f|      }t        j                  t        j
                  || j                   d d       | t        j                  |             S rs  r   r   _triboolrr  	broadcastr  r  r   r  r   Mmasks         r6   r<   r<   
  sZ    WW(1a	$A	"$	CMM$5q#:N:Nq:Q	RRr8   c                    | j                   ^ }}}t        j                  t        ||f|dz
        }t        j                  t        j
                  || j                   d d       t        j                  |       |       S )Nr   r   r  r  s         r6   rf  rf  
  s_    WW(1a	$AA	&$	CMM$5s7K7KA7NPQ	RRr8   c                    t        j                  d| j                  d         }t        j                  g | j                  | j                  d   d| j                        j
                  d||f   j                  |       S )z%Construct a (batched) diagonal matrixrW  r   r   .)r   r?  r   rM  r   r  r  r   r  s     r6   r  r  
  sd    	hhw$!	)AGG)QWWR[)1agg	6	9	9#q!)	D	H	H	KKr8   c                    t        j                  dt        | j                  d   | j                  d               }| d||f   S )z*Extract the diagonal from a batched matrixrW  r   r   .)r   r?  rN  r   r   s     r6   r  r  
  s9    	hhwAGGBK56!	
319r8   c           
         | j                   t        |      k  sJ t        j                  | |t	        t        |      | j                   z
  t        |                  S r   )r   r   r   r  r   )r?   r   s     r6   r  r  
  sE    	
3u:					ac%j166.A3u:(N	OOr8   c                   t        j                  |j                  t        j                        r9t        j                  | t        j                  t        j                  dz  z   |      S t        j                  | t        j                  |      S )Nr  )r   r   r   r   r   r   r[  r  )rK  r  s     r6   _nan_like_hlor  
  s[    tzz2#5#56sBFFRVVb[$8$??sBFFD11r8   c           	     V   t        t        j                  t        |j                        t        |j                        t        |j                                    }t        j
                  |||      }t        j                  | |||f|||f||      \  }}}t        j                  |||      S )z:Wrapper around XLA `Select` that broadcasts its arguments.)
r  r   r  r   r   broadcast_shardingsr   multi_broadcast_in_dimr%   rr  )	rK  which
which_avalr?   x_avalrt  y_aval
out_shapesr  s	            r6   _broadcasting_select_hlor  
  s    C((JuV\\2E&,,4GI J*((VVD,++C%A-7,H,6F+%A 
E1a	  r8   c                   t        |j                        t        |      z
  }t        |d|z  z   t        j                        }t        | t        j                  | ||t        t        |                  |||t        | |      |      S )Nr  r  )
r   r   r   r   bool_r  r   r  r   r  )rK  r   rd  r?   r
  num_bcast_dimsselect_avals          r6   r\  r\  
  sw    v||$s:6.J)>>I+	!	
C[16s:1GIsF+V
5 5r8   c                    t         j                  j                  t         j                  j	                  d      | j
                        S N   )r#   IntegerAttrrJ  IntegerTypeget_unsignedr5   )r  s    r6   r  r  
  s,    			BNN77:AGG	DDr8   c                    t         j                  j                  t         j                  j	                  d      t        |             S r  )r#   r  rJ  r  r  r  )r  s    r6   r#  r#  
  s,    			BNN77:CF	CCr8   c                (    t        | rd      S d      S )NrD  r  r"  rS  s    r6   r  r  
      	9C	..#	..r8   c                (    t        | rd      S d      S )NrD  r  r"  r   s    r6   rY  rY  
  s    	5C	**c	**r8   c                8    t        | r|rd      S d      S d      S )NCTr   r"  )r   	conjugates     r6   r  r  
  s!    	9IS	FF3	FF#	FFr8   c                (    t        | rd      S d      S )Nr  r   r"  )	unit_diags    r6   r  r  
  r  r8   c                J    | dz
  | dz
  ft        t        | dz
  dd            z   S )NrX   r   r  r   )r   r   )dims    r6   _column_major_matrix_layoutr$  
  s,    
'37	eE#'2r$:;	;;r8   c                     t        |j                        |z
  }|r|dk\  rdnd}|dj                   fdt        |      D              z   S )Nr   z...   c              3  4   K   | ]  }t                y wr   )rr  )r   r  letterss     r6   r   z%_sdy_rule_for_aval.<locals>.<genexpr>
  s     ;Q4=;s   )r   r   joinr   )r)  num_batch_dimsr  r   r   s   `    r6   _sdy_rule_for_avalr,  
  sB    	$**o&!#Q6B&	#((;%(;;	;;r8   c           
     v    t        t        j                        dj                   fd|D              }dj                   fd|D              }t	        | d|       }t        ||D cg c]  }t        j                  |       c}|D cg c]  }t        j                  |       c}      S c c}w c c}w )Nz, c              3  8   K   | ]  }t        |        y wr   r,  r   rs   r)  r+  s     r6   r   z+_build_sdy_sharding_rule.<locals>.<genexpr>
  s"      H9:.!4H   c              3  8   K   | ]  }t        |        y wr   r/  r0  s     r6   r   z+_build_sdy_sharding_rule.<locals>.<genexpr>
  s"      I9:.!4Ir1  z -> )iterstringascii_lettersr*  r   r   r   r2  )r+  rU  rP  lhsr^  sdy_sharding_rulers   r)  s   `      @r6   _build_sdy_sharding_ruler8  
  s    %%&'		 H>FH 	H#		 I>GI 	I#.#d3%/@A	"(011tA1(121tA2
4 412s   ,B1
B6c                *      fd}|S )Nc                   | j                   n}| j                  n}t        d g ||D              }xr | }t        d |D              }| j	                  ||      } |D cg c]<  }r6t        |j                        |k(  rt        t        |j                              nd > }	}|D cg c]<  }r6t        |j                        |k(  rt        t        |j                              nd > }
}|z
  }t        j                  dt        |      i      }|r/d|i}t        j                  j                  rt        |||      |d<   nd }t        j                   |	|
|      } || g|i |S c c}w c c}w )Nc              3  H   K   | ]  }t        |j                           y wr   r  )r   r  s     r6   r   z5_linalg_ffi_lowering.<locals>.rule.<locals>.<genexpr>
  s$      Q.2djj))Qr  c              3  F   K   | ]  }t        |j                          y wr   )r   r   )r   rn   s     r6   r   z5_linalg_ffi_lowering.<locals>.rule.<locals>.<genexpr>
  s     7s177|7s   !)rU  rP  r+  zmhlo.frontend_attributeszsdy.sharding_rule)operand_layoutsresult_layoutsrQ  extra_attributes)rU  rP  r  r  r  r   r   r$  r   ir_attributestrr   use_shardy_partitionerr5   r8  r   r  )rK  r  r  	avals_in_
avals_out_has_dynamic_shapebatch_partitionable_max_num_dimsr  r=  r>  r+  frontend_attrsr?  rb  rU  rP  r  r  r  rQ  r`  s                  r6   rb  z"_linalg_ffi_lowering.<locals>.rule
  s    ( 0hI"+"3J  Q6O	6OJ6OQ Q.H7H3H7Y77L
++yJ+
?C   C

O|; 	$C

O4AE	FO     C

O|; 	$C

O4AE	F N   "$66N&&(8#n:M'NON4nE		&	&	,	,0HIz13,- K+93I-=?D %d%f%%+ s   )AE10AE6rR   )r`  rU  rP  rQ  r  r  r  rb  s   ``````` r6   rX  rX  
  s     &  &B 
+r8   )r?   r'   r9   r  r  r'   )rD   r(   rE   r(   r  r'   )r?   r(   rS   r  rT   r  rU   zEigImplementation | NonerV   bool | Noner  zlist[Array])r?   r'   rg   r  r9   r  rh   r  ri   tuple[int, int] | NonerU   zEighImplementation | Noner  r  )rs   r(   r  r  )rs   r(   rw   r(   r  r'   )r?   r(   r  tuple[Array, Array, Array])r   r(   r~   r  r  r'   )
r?   r(   r   Literal[False]r   r  rV   rI  r  r  )
r?   r(   r   Literal[True]r   r  rV   rI  r  rK  )
r?   r(   r   r  r   r  rV   rI  r  z0tuple[Array, Array] | tuple[Array, Array, Array])
r?   r(   r   r  r   r  r   zCallable[..., Any] | Noner  r  )r?   r(   r   r  r   rM  ri   rJ  rl   SvdAlgorithm | Noner  rK  )r?   r(   r   r  r   rL  ri   rJ  rl   rN  r  r'   )r?   r(   r   r  r   r  ri   rJ  rl   rN  r  z"Array | tuple[Array, Array, Array])
r   r(   r   r(   r   r  r   r  r   r  )rs   r(   r   r(   r   r  rg   r  r   r  r   r  r   r  r  r'   )rs   r(   rg   r  r  z!tuple[Array, Array, Array, Array])
r   r'   r   r'   r   r'   r   r'   r  r'   )r   )FTT)r   rA  )   r  )
r|   r(   rh  r(   r   r(   r  r  r  r'   )
r|   r'   rh  r'   r   r'   r  r  r  r'   )rs   r(   r  r(   rV   rI  r  rK  )rs   r'   r   r'   r  r'   )rs   r'   r   r'   )r?   r'   r  r'   )r  r'   r   r  r  r'   )r   r'   r  r'   )r?   r'   r   tuple[int, ...]r  r'   )rK  zmlir.LoweringRuleContextr  ir.Value)r  rQ  )r   r  r  r  )r!  r  )r#  r  r  rP  )NNNTrX   T)
__future__r   collections.abcr   enum	functoolsr   r  r4  typingr   r   r   rZ   numpyr   jax._srcr	   r
   r   r   r   r   jax._src.corer   r   r   *jax._src.custom_partitioning_sharding_ruler   r   r   jax._src.interpretersr   r   r   jax._src.laxr   r   r   r#  jax._src.lax.laxr   r   r   jax._src.libr   r   r    r!   r"   jax._src.lib.mlirr#   jax._src.lib.mlir.dialectsr$   r%   jax._src.partition_specr&   r  jax._src.typingr'   r(   r7   r@   rF   EnumrH   r^   r`   rp   rt   rx   r|   r   ra   r   r   r   r   r   r   r   r   r   r  r  r  r.  input_dtypestandard_linalg_primitiver3  rG  rL  re  rh  r=   primitive_jvpsr   rm  r  r  rC   r  r  r  r  r  r  r  r  r]   r  r  r	  r  rm   r  r  r$  rr   r)  r5  r;  rv   rP  rX  r\  r`  rc  rq  rx  r  r  rz   r  r  jitr  r  r  r  r  r  rW  r   r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r   r  r  r  r   r  r  rg  r  r  r$  r.  r)  r   rE  rH  rR  r   rW  rY  rc  ri  rz  r  r   r  r  r  r  r  r  r   defjvp2primitive_transposesr(  r  r  r  r   r  r  r  r  r  r  r  r  r  r   r  r  r  r  r  r;   r<   rf  r  r  r  r  r  r\  r  r#  r  rY  r  r  r$  r,  r8  rX  rR   r8   r6   <module>rj     sV   # $     ) )         I I9  $ * & %  + 3 3 & # # #    + * 6 ,	; Z ( Z ( Z ( V $
 48 #<4&		  '+'+/3!O3O3  $O3 !%	O3
 -O3 O3 O3d  !!.20488 8 	8
 8 ,8 .8 8v$-"B1" 
HL#(; 
 
GK#(B 
 
).d#; 

 */d $33;3r #'15!'!'  !' 	!'
 /!' !'H499  
 .2%)  	
 , #   
 
 .2%)  	
 , #  
 
 .2%)  	
 , # ( 
 .2%)(( ( 	(
 ,( #( ((^ #&&& 	&
 & &Z 666 	6
 6 6 6 6 6t "&99 9&980:  %dEB  .E#<J@ @D"&B $$4cooF (8
KK 'h$ 4jB
 2  *    z#5 6   z#9E J *&<.>@@/ .Vf9;L    wA4H   DNN*UCEQ*#J:pd@ 	fx')4%	 (  %    u/% @ %!24D E
	---`'T 
v(*D2BF
 +  &  &"8 9
.*  Vh.0$8N4)   |%= N
 
>A
 2h)*F#%:<    ,.K L @B#FL.,
?"JR:& Vh&($ '  $    t^T^^JN O   t2U C $ 4 5 .. %.!( 	"( "(P"*!H 7hhZM4=     DNN4uMO  F&* )+T3Dg   w 4 5 '#: ; $( ,F*4"& )40&wEK '#: ;B
P2 D Vh&($ &  $    t^T^^L9 :
??+Z )+T3Dg   w 3e D
  !  + 2& ,2&h$$ D DLfR 	fx')4%	 )  %  %!6 7
:
Q 0Vf;    +T2VE   DNN,uEG AF 
&B<8A" RXXbjj)8288BJJ+?RXXbll+XRXXbmm-DF D@ & 6H#4fx6G"H
(*<>  

'FH /O  * +2Q  . /   )+E F   )+F %').2 !fx/14]TC -)F G
	!
H
).*.( 0h)6H+<fx>OP/1DF  *E  % &/P  + ,3S  / 0   #   +E   +F   *NDNNU-4 5
eIJ0H :S
S
L

P2!5ED/+G/<<

4 @DCGCG'r8   