
    SpfK                        d dl mZ d dlmZ d dlZd dlZd dlmc mZ d dl	mc m
c mZ d dlZddlmZ d dlmZ ddlmZmZmZmZmZmZmZ 	 ddlmZ n2# e$ r* d	D ]$Z	  ej         e d
          Z n# e$ r dZY !w xY wY nw xY wer: ej!                    "                                D ]\  Z#Z$ ej%        e#e$d           dD ]&Z	  ej         e dd          Z& n# e$ r dZ&Y #w xY we&r?e&!                                "                                D ]\  Z#Z$ ej%        e#e$d           	 ddl'mZ( n.# e$ r& dD ] Z)	  ej         e) d
          Z( n	#  dZ(Y xY wY nw xY we(r: e(j!                    "                                D ]\  Z#Z$ ej%        e#e$d           dD ]&Z)	  ej         e) dd          Z* n# e$ r dZ*Y #w xY we*r?e*!                                "                                D ]\  Z#Z$ ej%        e#e$d           d Z+d Z, ee,dee&          Z- ee,de(e*          Z.d Z/ ee/de&          Z0 ee/de*          Z1d Z2 ee2de          Z3 ee2de(          Z4d Z5 ee5de&          Z6d Z7 ee7de&          Z8 ee7de*          Z9ddde:edf         fd Z; ee;de&d!          Z< ee;de*d!          Z=	 d$d"Z> ee>de&d!          Z? ee>de*d          Z@d# ZA eeAde&          ZB eeAde*          ZCdS )%    )Sequence)partialN   )GpuLibNotLinkedError)
xla_client)DimensionSizeShapeTypePairmk_result_types_and_shapescustom_callensure_hlo_s32hlo_s32dense_int_array)_blas)jax_cuda12_pluginz._blasCUDA)platform)z.cudar   z._solverjaxlib)package)jax_rocm60_pluginROCM)z.rocmr   c                 4    t          j        |           j        S )z'Returns the real equivalent of 'dtype'.)npfinfodtype)r   s    Q/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jaxlib/gpu_solver.py
_real_typer   [   s    	%	    c                    t          j        |j                  }|j        }t	          |          dk    sJ |dd         \  }}t          |dd                   }	t	          |	          }
t          j        |	          }|st                      |dk    r||k    ry||z  dk    rp|	                    t          j        |          ||          \  }}t           j                            |gt           j                            d                    }|  d}nX|                    t          j        |          |||          \  }}t           j                            |g|j                  }|  d}|
|
dz   ft          t#          |
dz
  d	d	                    z   }t           j                            d
          }t%          ||j        t           j                            |	t'          ||          fz   |          t           j                            |	|          |g|g||g|t          t#          |
d	d	                    t          t#          |
dz
  d	d	                    dggddi          j        }|dd         S )zLU decomposition.   Nr         blas_getrf_batchedsolver_getrf    r   result_typesoperandsbackend_configoperand_layoutsresult_layoutsoperand_output_aliases   )irRankedTensorTypetypeshapelentuplemathprodr   build_getrf_batched_descriptorr   r   getIntegerTypeget_signlessbuild_getrf_descriptorelement_typeranger   minresults)r   gpu_blas
gpu_solverr   aa_typedimsmn
batch_dimsnum_bdbatchlworkopaque	workspacekernellayouti32_typeouts                      r   
_getrf_hlorQ   `   se   qv&&&	$	Ta	bcc$!QT#2#Y*z??&
)J

%	 !

 
  
QYY166a5jC//;;huooua! !ME6#''1L1LQ1O1OPPI,,,FF55
1& &ME6#''1DEEI&&&FFQJ%fqj"b(A(A"B"BB&^((,,(	

c!Qii\ 98DD

H55	 sheFB##$$eFQJB''((	
	  V#	% 	% 	%" &-# $ 
RaR.r   cuhipc                    t          j        |j                  }|j        }t	          |          dk    sJ |dd         \  }}t          |dd                   }t	          |          }	t          j        |          }
|                    t          j
        |          |
||          \  }}|	|	dz   ft          t          |	dz
  dd                    z   }t           j                            d          }t          |  d|j        t           j                            |t!          ||          fz   |j                  t           j                            ||          t           j                            |g|j                  g|g||g|t          t          |	dd                    t          t          |	dz
  dd                    dggddi	          j        }|dd
         S )zQR decomposition.r   r    Nr   r%   r&   solver_geqrfr   r'   r.   )r/   r0   r1   r2   r3   r4   r5   r6   build_geqrf_descriptorr   r   r=   r9   r:   r   r8   r>   r<   r?   )r   rA   r   rB   rC   rD   rE   rF   rG   rH   rI   rJ   rK   rN   rO   rP   s                   r   
_geqrf_hlorW      s   qv&&&	$	Ta	bcc$!QT#2#Y*z??&
)J

%33huooua$ $-% FQJ%fqj"b(A(A"B"BB&^((,,(		

c!Qii\ 96;NOO

H55
)<==	 sheFB##$$eFQJB''((	
	  V#	% 	% 	%" &-# $ 
RaR.r   c                    t          j        |j                  }|j        }t	          |          dk    sJ |dd         \  }}t          |dd                   }t	          |          }	t          j        |          }
|st                      |	                    t          j        |          |
||          \  }}|	|	dz   ft          t          |	dz
  dd                    z   }t          |  d|j        t           j                            |t          ||          fz   |j                  t           j                            |gt           j                            d                    t           j                            |gt           j                            d                    g|g||g|t          t          |	dd                    dgdggddi	          j        }|dd         S )
zBatched QR decomposition.r   r    Nr   r%   blas_geqrf_batchedr"   r   r'   )r/   r0   r1   r2   r3   r4   r5   r6   r   build_geqrf_batched_descriptorr   r   r=   r   r8   r>   r<   r9   r:   r?   )r   r@   r   rB   rC   rD   rE   rF   rG   rH   rI   rJ   rK   rN   rP   s                  r   _geqrf_batched_hlor[      s   qv&&&	$	Ta	bcc$!QT#2#Y*z??&
)J

%	 !

 
  99huooua$ $-% FQJ%fqj"b(A(A"B"BB&	%%%	

c!Qii\ 96;NOO
)D)DQ)G)GHH
)D)DQ)G)GHH	 sheFB##$$	
	
	  V#	 	 	$ % & 
RaR.r   c	                 H   t          j        |j                  }	t          j        |j                  }
|	j        d         }|
j        d         }|                    t          j        |          ||||          }t          |  d|j        g||||g|g ddg          j        }|S )z.Sparse solver via QR decomposition. CUDA only.r   solver_csrlsvqr)r   r^   r^   r^   r^   )r(   r)   r*   r+   r,   )	r/   r0   r1   r2   build_csrlsvqr_descriptorr   r   r   r?   )r   rA   r   dataindicesindptrbtolreorderb_type	data_typerF   nnzrK   rP   s                  r   _csrlsvqr_hlori      s     qv&&&!$),,)l1o!#//huooq#w & 		"""F8fgq)...V	 	 	   
*r   c                    t          j        |j                  }|j        }t	          |          dk    sJ |dd         \  }}t          |dd                   }	t	          |	          }
t          j        |	          }t          j        |j                  j        }|dd         |dd         k    sJ |d         }|                    t          j
        |          ||||          \  }}|
|
dz   ft          t          |
dz
  dd                    z   }t           j                            d          }t          |  d|j        t           j                            |	|          t           j                            |g|j                  g||g||t          t          |
dd                    g|t          t          |
dz
  dd                    dggddi	          j        }|dd         S )
z.Product of elementary Householder reflections.r   r    Nr%   r   r&   solver_orgqrr   r'   )r/   r0   r1   r2   r3   r4   r5   r6   build_orgqr_descriptorr   r   r=   r9   r:   r   r8   r<   r?   )r   rA   r   rB   taurC   rD   rE   rF   rG   rH   rI   tau_dimskrJ   rK   rN   rO   rP   s                      r   
_orgqr_hlorp      s   qv&&&	$	Ta	bcc$!QT#2#Y*z??&
)J

% **0(	#2#$ss)	#	#	#	#rl!33huoouaA' '-% FQJ%fqj"b(A(A"B"BB&^((,,(		

H55
)<==
 3x

fb"%%
&
&
 	eFQJB''((	

  V%	% 	% 	%$ &-% & 
RaR.r   F)lowera_shape_vals.c                   t          j        |j                  }t          |          dk    sJ |dd         \  }}	t          |          t          u rt          |	          t          u r||	k    s
J |            |dd         }
t          |
          }||dz   ft          t          |dz
  dd                    z   }t          d |
D                       }|rd}nt          j	        |
          }|r;|	dk    r5|s3|  d}|
                    t          j        |          |||	          \  }}n:|  d	}|                    t          j        |          |||	          \  }}|d
k    sJ t           j                            |j                  rt          j        |j                  j        }n|j        }t           j                            d          }|g}|g}|r`t'          d          }|
D ]$}t)          j        |t-          |                    }%|                    |           |                    d           ||j        f|
|	fz   |f|
|f|g|j        fg}t1          |          \  }}t3          ||||||t          t          |dd                    t          t          |dz
  dd                    d
ggd
d
i|          j        }|dd         S )z)Symmetric (Hermitian) eigendecomposition.r   r    Nr   r%   c              3   F   K   | ]}t          |          t          k    V  d S N)r1   int).0ds     r   	<genexpr>z_syevd_hlo.<locals>.<genexpr>3  s,      CCa477c>CCCCCCr   r&   solver_syevjsolver_syevdr    )r(   r)   r*   r+   r,   r-   result_shapesr.   )r/   r0   r1   r3   rv   r4   r=   anyr5   r6   build_syevj_descriptorr   r   build_syevd_descriptorComplexType
isinstancer<   r9   r:   r   hlomultiplyr   appendr
   r   r?   )r   rA   have_jacobi_solverr   rB   rr   rq   rC   rE   rF   batch_dims_valsrH   rN   dynamic_batch_dims	batch_intrM   rJ   rK   eigvals_typerO   r)   r+   batch_size_valb_vshape_type_pairsr(   r}   rP   s                               r   
_syevd_hlor   '  s    qv&&&	\		a					bcc	$!Q	aCDGGsNNqAvvv|vvv "%/&FQJ%fqj"b(A(A"B"BB&CC?CCCCC +II	/**I 
AGG,>G &&&F55
	1. .ME66 &&&F55
	1. .ME619999^v233 '>&"566CLL&L^((,,(S(H/ QZZN I I|NN34G4GHHnnOON###2 V()!|,!w#$	/&
 !;;K L L,%

fb"%%
&
&
fqj"b))
*
*#	  V!	# 	# 	# $+  
RaR.r   Tc                    t          j        |j                  }|j        }t	          |          dk    sJ |dd         \  }	}
t          |dd                   }t	          |          }t          j        |          }t           j        	                    |j
                  rt          j        |j
                  j
        }n|j
        }t          t          |dz
  dd                    }|ft          t          |dz
  dd                    z   }t           j                            d          }|r|	dk    r|
dk    r| o|dk    p|	dk    p|
dk    }|                    t          j        |          ||	|
||rdnd          \  }}t#          |	|
          }||dz   ft          t          |dz
  dd                    z   }t%          |  d	|j        t           j                            |t#          |	|
          fz   |          t           j                            ||	|r|n|	fz   |j
                  t           j                            ||
|r|n|
fz   |j
                  t           j                            ||          t           j                            |g|j
                  g|g||g|||||dggddi
          j        \  }}}}}}t+          j        |t/          t          j        t          t          |                    |dz   |fz                                 }t          j        |t          j                  rKt+          j        t+          j        |          t+          j        t+          j        |                              }|su|srt+          j        |t/          t          j         t	          |          gt          j!                            t/          t          j        ||	t#          |	|
          fz                       t/          t          j"        t	          |          gt          j!                                      }t+          j        |t/          t          j         t	          |          gt          j!                            t/          t          j        |t#          |	|
          |
fz                       t/          t          j"        t	          |          gt          j!                                      }n|	|
k     r]|#                    t          j        |          ||
|	||          \  }}|r|
n|	}|dz   |ft          t          |dz
  dd                    z   }t%          |  d|j        t           j                            |t#          |	|
          fz   |          t           j                            |||
fz   |j
                  t           j                            ||	|	fz   |j
                  t           j                            ||          t           j                            |g|j
                  g|g||g|||||dggddi
          j        \  }}}}}}n[|#                    t          j        |          ||	|
||          \  }}|r|	n|
}||dz   ft          t          |dz
  dd                    z   }t%          |  d|j        t           j                            |t#          |	|
          fz   |          t           j                            ||	|fz   |j
                  t           j                            ||
|
fz   |j
                  t           j                            ||          t           j                            |g|j
                  g|g||g|||||dggddi
          j        \  }}}}}}||||fS )zSingular value decomposition.r   r    Nr   r%   r&   i   r   solver_gesvdjr'   solver_gesvd)$r/   r0   r1   r2   r3   r4   r5   r6   r   r   r<   r=   r9   r:   build_gesvdj_descriptorr   r   r>   r   r8   r?   r   	transposer   array
issubdtypecomplexfloatingcomplexrealnegateimagslicezerosint64onesbuild_gesvd_descriptor)r   rA   r   r   rB   full_matrices
compute_uvrC   rD   rE   rF   rG   rH   rc   singular_vals_typescalar_layoutvector_layoutrO   econrJ   rK   ro   matrix_layout_suvinfovts                                r   
_gesvd_hlor   n  st    qv&&&	$	Ta	bcc$!QT#2#Y*z??&
i
!^v233 -(;<<I,fqj"b1122-)eE&1*b"$=$=>>>-^((,,(  e/AII!t)) =!q&"<AF"<a"fD66
Aq*4.>aaQ@ @ME6Aq		AVaZ(5vz2r1J1J+K+KKM%"""
&


!
!*Aq		|";=O
P
P


!
!*3C11!/D"D"("57 7


!
!*3C11!/D"D"("57 7


!
!*h
7
7


!
!5'6+>
?
?	
 &C
 !"1v/' ' '. (// Aq!Qa0 
	uV}}!5!5!V8L!LMMNN
P 
PB 
}UB.// ?;sx||SZ%=%=>>b 
; 
;
)

"(CII;99
:
:
"(:C1II#>??
@
@
"'3t99+rx88
9
9	; ;a
 9

"(CII;99
:
:
"(:QA#>??
@
@
"'3t99+rx88
9
9	; ;b
 	1uu55
Aq*m= =ME6!AaZ(5vz2r1J1J+K+KKM&!!!
&


!
!*Aq		|";=O
P
P


!
!*1v"5v7J
K
K


!
!*1v"5v7J
K
K


!
!*h
7
7


!
!5'6+>
?
?
 &




#
 !"1v+' ' '* (/+ Aq"aqq. 55
Aq*m= =ME6!AVaZ(5vz2r1J1J+K+KKM&!!!
&


!
!*Aq		|";=O
P
P


!
!*1v"5v7J
K
K


!
!*1v"5v7J
K
K


!
!*h
7
7


!
!5'6+>
?
?
 &




#
 !"1v+' ' '* (/+ Aq!Rq, 
Ar4r   c                   t          j        |j                  }|j        }t	          |          dk    sJ |dd         \  }}||k    sJ ||f            t          |dd                   }	t	          |	          }
d}|	D ]}||z  }|                    ||||          \  }}t          j        |t          j	                  r|j
        }np|t          j        k    rt           j                                        }nA|t          j        k    rt           j                                        }nt!          d|           |
|
dz   ft          t#          |
dz
  dd                    z   }t           j                            d          }t)          |  d|j        t           j                            |	|fz   |          t           j                            |	|dz
  fz   |          t           j                            |	|dz
  fz   |j
                  t           j                            |	|          t           j                            |g|j
                  g|g||g||
ft          t#          |
dz
  dd                    z   |
ft          t#          |
dz
  dd                    z   |
ft          t#          |
dz
  dd                    z   t          t#          |
dz
  dd                    d	ggd	d	i
          j        \  }}}}}}d d }|sK| dk    rD|dk    r=dt	          |	          z  dz   }|	dz   }t-          j        | ||           ||           |dgt	          |          z                      }t           j                            |	dz   |          }t-          j        || |t#          t	          |          dz
                                }t-          j        t           j                            |j        |j
                  |          }t          fddt	          |	          z  dz   D                       }t-          j        |||          }|||||fS )zGsytrd: Reduction of a symmetric (Hermitian) matrix to tridiagonal form.r   r    Nr   zUnsupported dtype r%   r&   solver_sytrdr   r'   c                 z    t           j                            t          j        | t          j                            S ru   )r/   DenseIntElementsAttrr8   r   asarrayr   xss    r   <lambda>z_sytrd_hlo.<locals>.<lambda>&  s%    r.222:b"(3K3KLL r   c                 Z    t          t          j        | t          j                            S ru   )r   r   r   r   r   s    r   r   z_sytrd_hlo.<locals>.<lambda>'  s    /"*R*B*BCC r   rR   r^   )r   )r   r   c              3   R   K   | ]!}t          j         |                    V  "d S ru   )r   constant)rw   iintattrs     r   ry   z_sytrd_hlo.<locals>.<genexpr>2  sQ       @ @ L,, @ @ @ @ @ @r   )r   r   )r/   r0   r1   r2   r3   r4   build_sytrd_descriptorr   r   floatingr<   	complex64F32Typer8   
complex128F64TypeNotImplementedErrorr=   r9   r:   r   r?   r   r   broadcast_in_dimconvertdynamic_update_slice)r   rA   r   rB   rq   rC   rD   rE   rF   rG   rH   rc   rx   rJ   rK   	diag_typerN   rO   etausr   r   
intarrattrstartendr   s_typeoffsetsr   s                               @r   
_sytrd_hlor     sb   qv&&&	$	Ta	bcc$!Q	
a!QT#2#Y*z??&!  aFAA33E5!QGG-%]5"+&& <#II
  II
  II
:5::
;
;;FQJ%fqj"b(A(A"B"BB&^((,,(&		

aT 19==

a!eX 5yAA

a!eX 5v7JKK

H55
)<== sh		E%
B33444		E%
B33444		E%
B33444eFQJB''((	
  V+  , - !Q4q6 ML'CC*	 08t##A3z??"T)E
t
C		::ejjoozz1#E

:J/K/K	M 	MA $$Z&%8)DDFVQ

5TQ3G3G(H(HIIA
f.ABBA	G 	GA @ @ @ @"S__4v=@ @ @ @ @G Aw//A	
Aq$	r   )TT)Dcollections.abcr   	functoolsr   	importlibr5   jaxlib.mlir.irmlirr/   jaxlib.mlir.dialects.stablehlodialects	stablehlor   numpyr   gpu_common_utilsr   r   r   hlo_helpersr   r	   r
   r   r   r   r   cudar   _cublasImportErrorcuda_module_nameimport_moduleregistrationsitems_name_valueregister_custom_call_target	_cusolverrocm_hipblasrocm_module_name
_hipsolverr   rQ   
cuda_getrf
rocm_getrfrW   
cuda_geqrf
rocm_geqrfr[   cuda_geqrf_batchedrocm_geqrf_batchedri   cuda_csrlsvqrrp   
cuda_orgqr
rocm_orgqrr4   r   
cuda_syevd
rocm_syevdr   
cuda_gesvd
rocm_gesvdr   
cuda_sytrd
rocm_sytrdr|   r   r   <module>r      sQ   % $ $ $ $ $                     , , , , , , , , , , , ,     2 2 2 2 2 2      ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ;	$$$$$$$   /  '	'+;(C(C(CDDg e    ggg	  K,w,..4466 K KmeV*J*5&6JJJJJ6 
 

'	'%%%x  I 
E 
   III
  K ..006688 K KmeV*J*5&6JJJJJ	%%%%%%%   /  ((,<)D)D)DEEh ehhh  M-x-//5577 M MmeV,j,UFVLLLLL6 
 

((%%%x  J 
E 
   JJJ
  M!//117799 M MmeV,j,UFVLLLLL  
, , ,\ WZw	::
WZ*==
! ! !F WZy11
WZ
33
$ $ $L W/w?? W/AA   , tY77& & &P WZy11
WZ
33
 ?DA A A"=##56A A A AF WZy$77
WZ
D99
 /3~ ~ ~ ~@ WZy$77
WZ
E::
D D DL WZy11
WZ
33


s   A 	BA31B3A=:B<A==BBCC('C(-D4 4	E>EEEEE#F::GG