
    Spfi4                        d Z ddlZddlmZ ddlZddlmc mZ ddlZ	ddl
mZ ddlmZmZ 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 ]&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            eeoej                  Z eeoej                  Zd Zd Zd Z  ee de          Z! ee de          Z"d Z# ee#de          Z$ ee#de          Z%dddddZ& ee&de          Z' ee&de          Z(dddddZ) ee)de          Z* ee)de          Z+d Z, ee,de          Z- ee,de          Z.d Z/ ee/de          Z0 ee/de          Z1dddddZ2 ee2de          Z3 ee2de          Z4dddddZ5 ee5de          Z6 ee5de          Z7dddZ8 ee8de          Z9 ee8de          Z:dS )zD
cusparse wrappers for performing sparse matrix computations in JAX
    N)partial)
xla_client   )custom_callmk_result_types_and_shapes)z.cudajax_cuda12_pluginz._sparsejaxlib)packageCUDA)platform)z.rocmjax_rocm60_pluginROCMc                 6   t          j        | j                  }t          j        |j                  }t          j        |j                  }|j        \  }|j        |gk    sJ |j        |j        k    sJ |j        |d         dz   gk    sJ |j        |j        |fS )Nr   r   irRankedTensorTypetypeshapeelement_type)dataindicesindptrr   	data_typeindices_typeindptr_typennzs           Q/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jaxlib/gpu_sparse.py_validate_csr_hlor   ?   s    !$),,)$W\22,#FK00+	$#		u	$	$	$	$		!\%>	>	>	>	>		uQx!|n	,	,	,	,		!:C	??    c                 $   t          j        | j                  }t          j        |j                  }t          j        |j                  }|j        \  }|j        |gk    sJ |j        |j        k    sJ |j        |gk    sJ |j        |j        |fS )Nr   )r   rowcolr   row_typecol_typer   s          r   _validate_coo_hlor%   J   s    !$),,) **( **(	$#	C5	 	 	 	 		("7	7	7	7	7	C5	 	 	 	 		!6	;;r   c          
         t          ||||          \  }}	}
|\  }}|                    |||||
          \  }}t          |  dt          j                            ||          t          j                            |gt          j                            d                    g|||g|dggdz  ddgdgg          j        }|d         S )zCSR to dense matrix.sparse_csr_todense   r      r   result_typesoperandsbackend_configoperand_layoutsresult_layouts)	r   build_csr_todense_descriptorr   r   r   getIntegerTypeget_signlessresults)r   
gpu_sparser   r   r   r   
data_dtypeindex_dtyper   
index_typer   rowscolsbuffer_sizeopaqueouts                   r   _csr_todense_hlor>   V   s      1wNN)Z*$"??+tT30 0+v 		%%%


!
!%
3
3


!
!;-"$."="=a"@"@B B
 gv&seai!fqc]
	$ 
	$ 
	$ %,  
Q-r   cuhipc                *   t          j        |j                  }|j        \  }}	|                    ||||	|          \  }
}t          |  dt           j                            |g|j                  t           j                            |g|          t           j                            |dz   g|          t           j                            |
gt           j        	                    d                    g|g|ddggdggdz            j
        }|dd         S )	CSR from dense matrix.sparse_csr_fromdenser   r(   r      r*   Nr)   )r   r   r   r   build_csr_fromdense_descriptorr   r1   r   r2   r3   r4   )r   r5   matr   r7   r6   r8   mat_typer9   r:   r;   r<   r=   s                r   _csr_fromdense_hlorH   p   s      **(~*$"AA+tT30 0+v 		'''


!
!3%)>
?
?


!
!3%
4
4


!
!4!8*j
9
9


!
!;-"$."="=a"@"@B B	 u1vhcUQY	  	  	  !(  
RaR.r   F)	transposecompute_dtypecompute_typec          
         t          ||||          \  }}}|\  }}||
}|}	|                    |
|||||||          \  }}|r|n|}t          |  dt          j                            |g|	          t          j                            |gt          j                            d                    g||||g|dggdz  dggdz            j        }|d         S )zCSR matrix/vector multiply.Nsparse_csr_matvecr(   r   rD      r*   )	r   build_csr_matvec_descriptorr   r   r   r1   r2   r3   r4   )r   r5   r   r   r   xr   rI   rJ   rK   r6   r7   x_dtyper   r8   r   r9   r:   r;   r<   out_sizer=   s                         r   _csr_matvec_hlorS      s     1wNN)Z*$ML">>'=+
D#y" "+v (TTD(	$$$


!
!8*l
;
;


!
!;-"$."="=a"@"@B B
 gvq)seaicUQY
	  
	  
	  !(  
Q-r   c                    t          ||||          \  }}}|\  }}t          j        |j                  j        }|\  }}||}|}	|                    ||||
|||||	  	        \  }}|r|n|}t          |  dt          j                            ||g|	          t          j                            |gt          j        	                    d                    g||||g|dgdgdgddggddgdgg          j
        }|d         S )rB   Nsparse_csr_matmatr(   r   r   r*   )r   r   r   r   r   build_csr_matmat_descriptorr   r1   r2   r3   r4   )r   r5   r   r   r   Br   rI   rJ   rK   r7   r6   B_dtyper   r8   r   r9   r:   B_shape_Ccolsr;   r<   rR   r=   s                            r   _csr_matmat_hlor\      sI     1wNN)Z*$''-'(!UML">>'=+
D%i) )+v (TTD(	$$$


!
!8U"3\
B
B


!
!;-"$."="=a"@"@B B
 gvq)sQC!q!f-!fqc]
	$ 
	$ 
	$ %,  
Q-r   c          
         t          |||          \  }}	}
|\  }}|                    |||||
          \  }}t          |  dt          j                            ||          t          j                            |gt          j                            d                    g|||g|dggdz  ddgdgg          j        }|d         S )zCOO to dense matrix.sparse_coo_todenser(   r   r)   r   r*   )	r%   build_coo_todense_descriptorr   r   r   r1   r2   r3   r4   )r   r5   r   r!   r"   r   r6   r7   r   rZ   r   r9   r:   r;   r<   r=   s                   r   _coo_todense_hlor`      s     (c377)Q*$"??+tT30 0+v 		%%%


!
!%
3
3


!
!;-"$."="=a"@"@B B
 c3seai!fqc]
	$ 
	$ 
	$ %,  
Q-r   c                $   t          j        |j                  }|j        \  }}	|                    ||||	|          \  }
}t          |  dt           j                            |g|j                  t           j                            |g|          t           j                            |g|          t           j                            |
gt           j        	                    d                    g|g|ddggdggdz            j
        }|dd         S )	COO from dense matrix.sparse_coo_fromdenser(   r   r   rD   r*   Nr)   )r   r   r   r   build_coo_fromdense_descriptorr   r1   r   r2   r3   r4   )r   r5   rF   r   r6   r7   r8   rG   r9   r:   r;   r<   r=   s                r   _coo_fromdense_hlore      s     **(~*$"AA+tT30 0+v 		'''


!
!3%)>
?
?


!
!3%
4
4


!
!3%
4
4


!
!;-"$."="=a"@"@B B	 u1vhcUQY	  	  	  !(  
RaR.r   c          
         t          |||          \  }}}|\  }}||}|}	|                    ||||
||||          \  }}|r|n|}t          |  dt          j                            |g|	          t          j                            |gt          j                            d                    g||||g|dggdz  dggdz            j        }|d         S )zCOO matrix/vector multiply.Nsparse_coo_matvecr(   r   rD   rN   r*   )	r%   build_coo_matvec_descriptorr   r   r   r1   r2   r3   r4   )r   r5   r   r!   r"   rP   r   rI   rJ   rK   r7   r6   rQ   r   rZ   r   r9   r:   r;   r<   rR   r=   s                         r   _coo_matvec_hlori     s    (c377)Q*$ML">>'=+
D#y" "+v (TTD(	$$$


!
!8*l
;
;


!
!;-"$."="=a"@"@B B
 c3"seaicUQY
	  
	  
	  !(  
Q-r   c                   t          |||          \  }}}d}d}t          |          dk    r|\  }}n t          |          dk    rd}|\  }}}||z  }t          j        |j                  j        }|\  }}||}|}	|}|r|n|}||z  }|                    ||
||||||||||          \  }}|r|n|}|r
|||g}g d}n||g}ddg}t          |  d	t          j                            ||	          t          j                            |gt          j	        
                    d
                    g||||g|dgdgdgddgg|dgg          j        }|d         S )rb   Fr   rN   r)   TN)rN   r   r   r   sparse_coo_matmatr(   r*   )r%   lenr   r   r   r   build_coo_matmat_descriptorr   r1   r2   r3   r4   )r   r5   r   r!   r"   rW   r   rI   rJ   rK   rQ   r6   r7   r   rZ   r   is_batched_matmatbatch_countr9   r:   rY   r[   lhs_batch_strideB_rowsrhs_batch_strider;   r<   rR   	out_shape
out_layoutr=   s                                  r   _coo_matmat_hloru   '  s    (c377)Q+ZZ1__JD$$
5zzQ#Kt

C''-'(!UML &44$&un">>'=+
D%i6F +v (TTD( h.IJJ5!IQJ	$$$


!
!)\
:
:


!
!;-"$."="=a"@"@B B
 c3"sQC!q!f- 1#&
	( 
	( 
	( )0  
Q-r   )b_shape_valsc                
   t          |
          dk    sJ |
dd         }t          j        |          }t          |
          dz
  }|	t          j        k    }|r|                    |||          }n|                    |||          }||dz   ft          t          |dz
  dd                    z   }|ft          t          |dz
  dd                    z   }t          j
        |j                  }|||fz   |j        f|ft          j                            d          fg}t          |          \  }}t!          |  d|rdnd	z   |||||g|                    ||||          |gd
z  |gz   |dggd
di|          j        }|d         S )z2Calls `cusparse<t>gtsv2(dl, d, du, B, m, n, ldb)`.rN   Nr   r(   sparse_gtsv2_f32f64r)   r   )r+   r,   r-   r.   r/   operand_output_aliasesresult_shapes)rl   mathprodnpfloat32gtsv2_f32_buffer_sizegtsv2_f64_buffer_sizetupleranger   r   r   r   r2   r3   r   r   build_gtsv2_descriptorr4   )r   r5   dlddurW   mnldbtrv   batch_dim_vals
batch_sizenum_bdr{   r;   b_layoutd_layoutb_typeshape_type_pairsr+   r~   r=   s                          r   
_gtsv2_hlor   c  s    
\		a				$.y((*|q &	
bj# >221a==KK221a==Kfqj!E%
B*C*C$D$DD(YuVaZR88999(qv&&& a &"56~r~221556 !;;K L L,	   S$;EEe<Ar1~66z1aMMj1nz1_V!	# 	# 	# $+  
Q-r   );__doc__r   	functoolsr   	importlibjaxlib.mlir.irmlirr   numpyr   r	   r   hlo_helpersr   r   cuda_module_nameimport_module	_cusparseImportErrorregistrationsitems_name_valueregister_custom_call_targetrocm_module_name
_hipsparseboolsparse_supportedcuda_is_supportedrocm_is_supportedr   r%   r>   cuda_csr_todenserocm_csr_todenserH   cuda_csr_fromdenserocm_csr_fromdenserS   cuda_csr_matvecrocm_csr_matvecr\   cuda_csr_matmatrocm_csr_matmatr`   cuda_coo_todenserocm_coo_todensere   cuda_coo_fromdenserocm_coo_fromdenseri   cuda_coo_matvecrocm_coo_matvecru   cuda_coo_matmatrocm_coo_matmatr   
cuda_gtsv2
rocm_gtsv2 r   r   <module>r      s                                  @ @ @ @ @ @ @ @6 
 

'	'%%%x  I 
E 
   III
  K ..006688 K KmeV*J*5&6JJJJJ6 
 

((%%%x  J 
E 
   JJJ
  K!//117799 K KmeV*J*5&6JJJJJ DAy'ABB DC
(CDD 	@ 	@ 	@	< 	< 	<  , 7+T9== 7+UJ??   0 W/yAA W/
CC  $4d    : '/4;;'/5*== $4d    > '/4;;'/5*==  , 7+T9== 7+UJ??   0 W/yAA W/
CC  $4d    : '/4;;'/5*== $4d6 6 6 6 6p '/4;;'/5*== GK    B WZy11
WZ
33


s#   AAAB22B<;B<