
    `i                        d dl Z d dlZd dlZ	 d dlZdZn# e$ r dZY nw xY w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 d dlmZ  G d dej                  Zd Zd,dZd Zd ZdZdZej                            d          d             Zej                            d          d             Zd Zej                            d          d             Zej                            d          d             Z dZ!dZ"dZ#dZ$dZ%dZ&d Z'd!Z(d" Z)ej                            d          d-d$            Z*ej                            d          d%             Z+d& Z,ej                            d          d'             Z-d( Z.ej                            d          d)             Z/ej                            d          d*             Z0ej                            d          d+             Z1dS ).    NTF)_accelerator)cub)runtime)_base)_compressed)_csc)SparseEfficiencyWarning)_utilc                   *   e Zd ZdZdZd0dZd Zd Zd Zd Z	d	 Z
d
 Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd1dZd Zd Zd Zd Zd Zd1dZd Zd2dZd3dZd4d Zd4d!Z d4d"Z!d# Z"d4d$Z#d4d%Z$d4d&Z%d3d'Z&d( Z'd) Z(d* Z)d+ Z*d, Z+d- Z,d. Z-d/ Z.dS )5
csr_matrixa  Compressed Sparse Row matrix.

    This can be instantiated in several ways.

    ``csr_matrix(D)``
        ``D`` is a rank-2 :class:`cupy.ndarray`.
    ``csr_matrix(S)``
        ``S`` is another sparse matrix. It is equivalent to ``S.tocsr()``.
    ``csr_matrix((M, N), [dtype])``
        It constructs an empty matrix whose shape is ``(M, N)``. Default dtype
        is float64.
    ``csr_matrix((data, (row, col)))``
        All ``data``, ``row`` and ``col`` are one-dimenaional
        :class:`cupy.ndarray`.
    ``csr_matrix((data, indices, indptr))``
        All ``data``, ``indices`` and ``indptr`` are one-dimenaional
        :class:`cupy.ndarray`.

    Args:
        arg1: Arguments for the initializer.
        shape (tuple): Shape of a matrix. Its length must be two.
        dtype: Data type. It must be an argument of :class:`numpy.dtype`.
        copy (bool): If ``True``, copies of given arrays are always used.

    .. seealso::
        :class:`scipy.sparse.csr_matrix`

    csrNc                    t           st          d          | j                            |          }| j                            |          }| j                            |          }t          j                            |||f| j	                  S )a:  Returns a copy of the array on host memory.

        Args:
            stream (cupy.cuda.Stream): CUDA stream object. If it is given, the
                copy runs asynchronously. Otherwise, the copy is synchronous.

        Returns:
            scipy.sparse.csr_matrix: Copy of the array on host memory.

        zscipy is not availableshape)
_scipy_availableRuntimeErrordatagetindicesindptrscipysparser   _shape)selfstreamr   r   r   s        k/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/cupyx/scipy/sparse/_csr.pyr   zcsr_matrix.get8   s       	97888y}}V$$,""6**((|&&7F#4; ' 8 8 	8    c                 H    t          |          }|j        |j        |j        fS N)	dense2csrr   r   r   )r   xms      r   _convert_densezcsr_matrix._convert_denseK   s     aLLvqy!(**r   c                 
    ||fS r    )r   r!   ys      r   _swapzcsr_matrix._swapO   s    1vr   c                 $   ddl m} |                                  |                                }|                                 |                    d          r|j        }n$|                    d          r|j        }nt           || |||          S )Nr   cusparsecsrgeam2csrgeam)cupyxr*   sum_duplicatestocsrcheck_availabilityr+   r,   NotImplementedError)r   otheralphabetar*   r,   s         r   _add_sparsezcsr_matrix._add_sparseR   s    """"""&&z22 	&'GG((33 	&&GG%%wtUE4000r   c                 4   t          j        |          rt          j        || j                                      d          }t          j        |d                   rX|dk    r2t          t          j	        | j
        t          j                            S t          | j
        t          j                  S t          j        dt          j                  }t          j        dt          j                  }t          |||fd          }t          | ||          S t          j        |          r ||                                 |          S t%          |          r|                                  |                                 |d	v rt          | ||          S t)          j        d
t,                     |dk    rd}n|dk    rd}n|dk    rd}t          | ||          }t          j        |                                          }	t          |	          S t2          )Ndtype   r   _ne_)r9      )r9   r9   r   )r:   _lt__gt_z]Comparing sparse matrices using ==, <=, and >= is inefficient, try using !=, <, or > instead._eq__le_r=   _ge_r<   )r
   isscalarlikecupyasarrayr8   reshapenumpyisnanr   onesr   bool_zerosint32arange
binopt_csrisdensetodenseisspmatrix_csrr.   warningswarnr	   logical_nottoarrayr1   )
r   r2   opop_namer   r   r   opposite_op_nameresouts
             r   _comparisonzcsr_matrix._comparison`   s   e$$ 	#<TZ888@@CCD{47## Ef$$%di
%+&N&N&NOOO%djDDDDjU[999G[%+666Fgv6fEEEEdE7333]5!! 	#2dllnne,,,E"" 	#!!!  """222!$w777M>') ) ) &  #)  F""#)  F""#) T5*:;;C"3;;==11Cc??"!!r   c                 D    |                      |t          j        d          S )Nr>   )rY   operatoreqr   r2   s     r   __eq__zcsr_matrix.__eq__       x{F;;;r   c                 D    |                      |t          j        d          S )Nr:   )rY   r[   ner]   s     r   __ne__zcsr_matrix.__ne__   r_   r   c                 D    |                      |t          j        d          S )Nr<   )rY   r[   ltr]   s     r   __lt__zcsr_matrix.__lt__   r_   r   c                 D    |                      |t          j        d          S )Nr=   )rY   r[   gtr]   s     r   __gt__zcsr_matrix.__gt__   r_   r   c                 D    |                      |t          j        d          S )Nr?   )rY   r[   ler]   s     r   __le__zcsr_matrix.__le__   r_   r   c                 D    |                      |t          j        d          S )Nr@   )rY   r[   ger]   s     r   __ge__zcsr_matrix.__ge__   r_   r   c           
      	   ddl m} t          j        |          r1|                                  |                     | j        |z            S t          |          r|                                  |                                 |                    d          r|	                    | |          S |                    d          r|
                    | |          S |                    d          r|                    | |          S t          t          j        |          r|                                  |                                 |                    d          r)t          j        s|                    | |j        d          S |                    d          r>|                                }|                                 |	                    | |          S |                    d          r>|                                }|                                 |
                    | |          S t          t'          j        |          r| |                                z  S t'          j        |          r^|j        dk    r1|                                  |                     | j        |z            S |j        dk    r|                                  t          j        |          }| j        j        j        j        | j        j        | j        j        j        z  k    }|t;          j                    d	k     z  }t?          j                     D ]r}|t>          j!        k    r`t          j        sT|rR|j"        j#        rFt;          j$        | j%        d         | j%        d         | j&        | j        | j        | j'        |          c S s|                    d
          r)| j&        dk    r|(                    | |          r|j)        }nA|                    d          r|j*        }n$|                    d          r|j+        }nt           || |          S |j        dk    rs|                                  |                    d          r|j,        }n$|                    d          r|j-        }nt           || t          j        |                    S t]          d          t^          S )Nr   r)   spgemmcsrgemm2csrgemmT)transbr9   i*  csrmvExcsrmvspmvr;   csrmm2spmmzcould not interpret dimensions)0r-   r*   rB   isscalarr.   
_with_datar   rO   r0   rp   rq   rr   AssertionErrorr   isspmatrix_cscr   is_hipTr/   r   
isspmatrixrM   ndimasfortranarrayr   memsizer8   itemsizer   _get_cuda_build_versionr   get_routine_acceleratorsACCELERATOR_CUBflagsc_contiguousdevice_csrmvr   nnzr   csrmvExIsAlignedrt   ru   rv   rw   rx   
ValueErrorNotImplemented)r   r2   r*   bis_cub_safeacceleratorru   csrmms           r   __mul__zcsr_matrix.__mul__   s   """"""= L	"!!!??49u#4555E"" I	"!!!  """**844 %tU333,,Z88 %((u555,,Y77 %''e444$$ '' >	"!!!  """**955 %gn %''egd'CCC,,X66 	%KKMM  """tQ///,,Z88 %KKMM  """((q111$$e$$ .	"%++--'']5!! ,	"zQ##%%%ty5'8999q##%%%+E22  ${/38!%!1DK4E4N!N O  ; = = EF#/#H#J#J I IK#|'CCC$+N D + D050H D  #/ JqM4:a=$( It{DL% I  I I I I //	:: 	)tx!|| 11$>> @L %,EE0099 )$NEE0088 )$MEE((uT5)))q##%%%..x88 )$OEE0088 )$MEE((uT4#6u#=#=>>> !ABBB!!r   c                     t           r   r1   r]   s     r   __div__zcsr_matrix.__div__       !!r   c                     t           r   r   r]   s     r   __rdiv__zcsr_matrix.__rdiv__   r   r   c                    t          j        |          r^| j        }|t          j        k    rt          j        }t          j        ||          }t          j        ||          }t          | |          S t          j
        |          rt          j        |          }t          j        || j                  }t          | j        |j                   |                                 } t!                      |j        |j        |j        |j        d         |          |_        |S t)          j        |          rt          | j        |j        d           t          j        | j        |j                  }|j        dvrt          j        t          j        |          }|                                                     |d          }||                                z  S t4          S )z7Point-wise division by another matrix, vector or scalarr7   r9   F)allow_broadcastingFDcopy)r
   rA   r8   rE   float32float64rB   result_type
reciprocalmultiply_by_scalarrM   
atleast_2dbroadcast_tor   check_shape_for_pointwise_optocoo_cupy_divide_by_denser   rowcolr   r   promote_typescharrN   astyper   )r   r2   r8   dret
self_denses         r   __truediv__zcsr_matrix.__truediv__   s   e$$ 	0JE%% $UE22EU333A%dA...]5!! 	0OE**E%eTZ88E(U[AAA**,,C.,..#'37CIaL%A ACHJe$$ 	0 )U[<AC C C C'
EK@@Ez%%+EM5AA ..u5.AAJ//r   c                     t           S r   )r   r]   s     r   __rtruediv__zcsr_matrix.__rtruediv__  s    r   r   c           	      v   | j         \  }}t          |t          |d          z   |t          |d          z
            }|dk    rt          j        d| j                  S |                                  t          j        || j                  } t                      |||| j        | j	        | j
        |           |S )Nr   r7   )r   minmaxrB   emptyr8   r.   _cupy_csr_diagonalr   r   r   )r   krowscolsylenr&   s         r   diagonalzcsr_matrix.diagonal  s    Z
d4#a))#TC1II%566199:atz2222Jt4:...QdDIt{!\1	. 	. 	.r   c                     ddl m} |                    | d          }|j        | _        |j        | _        |j        | _        dS )zRemoves zero entories in place.r   r)   N)r-   r*   csr2csr_compressr   r   r   )r   r*   compresss      r   eliminate_zeroszcsr_matrix.eliminate_zeros  sI    """""",,T155M	'or   c                    t          j        |          rt          j        || j                  } ||          r| j        }|t
          j        k    rt
          j        }n|t
          j        k    rt
          j	        }t          j
        ||          }|                    |d          } ||                                 |          }t          |          S |                                   || j        |          }t          || j        | j        f| j        | j                  S t          j        |          rF|                                  t          j        |          } ||                                 |          S t+          |          r9|                                  |                                 t-          | ||          S t.          )Nr7   Fr   )r   r8   )r
   rA   rB   rC   r8   rE   r   r   	complex64
complex128r   r   rN   r   r.   r   r   r   r   rM   r   rO   rL   r1   )r   r2   cupy_oprU   dense_checkr8   	new_arraynew_datas           r   _maximum_minimumzcsr_matrix._maximum_minimum(  s   e$$ 	4Ldj999E{5!! F
 EM))!MEEeo--!,E(66U77#GDLLNNE::	!),,,##%%%"749e44!8T\4;"G(,
$*F F F F]5!! 	4!!!OE**E74<<>>5111E"" 	4!!!  """dE7333!!r   c                 H    |                      |t          j        dd           S )N	_maximum_c                     | dk    S Nr   r%   r!   s    r   <lambda>z$csr_matrix.maximum.<locals>.<lambda>I  
    q1u r   )r   rB   maximumr]   s     r   r   zcsr_matrix.maximumG  '    $$UDL+%4_6 6 	6r   c                 H    |                      |t          j        dd           S )N	_minimum_c                     | dk     S r   r%   r   s    r   r   z$csr_matrix.minimum.<locals>.<lambda>M  r   r   )r   rB   minimumr]   s     r   r   zcsr_matrix.minimumK  r   r   c                    t          j        |          rt          | |          S t          j        |          r8|                                  t          j        |          }t          | |          S t          |          r8|                                  |                                 t          | |          S d}t          |          )z=Point-wise multiplication by another matrix, vector or scalarz2expected scalar, dense matrix/vector or csr matrix)rB   ry   r   r
   rM   r.   r   multiply_by_denserO   multiply_by_csr	TypeError)r   r2   msgs      r   multiplyzcsr_matrix.multiplyO  s    = 	!%dE222]5!! 
	!!!!OE**E$T5111E"" 	!!!!  """"4///FCC.. r   c                    | j         \  }}t          d|           t          d|          }}t          ||z
  ||z
            }|dk    rt          d          |                    | j                  }|j        dk    rt          j        |f|| j                  }nt          ||j	                  }|d|         }t          j
        |||z   d          }	t          j        |dz   fd          }
t          j
        |dz   d          |
|||z   dz   <   ||
||z   dz   d<   ||                     |          d|         z  }| t          ||	|
f| j                   z   }|j        | _        |j        | _        |j        | _        dS )	z3Set diagonal or off-diagonal elements of the array.r   zk exceeds matrix dimensionsr7   Nir9   )r   r   )r   r   r   r   r   r8   r   rB   fullr   rK   rI   r   r   r   r   r   )r   valuesr   r   r   row_stcol_stx_lenx_data	x_indicesx_indptrr&   s               r   setdiagzcsr_matrix.setdiaga  s   Z
dQSAYYD6M4&=11A:::;;;tz**;!YxtzBBBFFv{++EFUF^FKcBBB	:taxk555*.+eAgS*I*I*IuQ&'$)a!$--!-$$VeV,,:vy(;4:NNNNF	yhr   c                 \    ddl m} | j        s|                    |            d| _        dS dS )zSorts the indices of this matrix *in place*.

        .. warning::
            Calling this function might synchronize the device.

        r   r)   TN)r-   r*   has_sorted_indicescsrsort)r   r*   s     r   sort_indiceszcsr_matrix.sort_indicesy  sM     	#"""""& 	+T"""&*D###	+ 	+r   c                    ddl m} |dn|                                }| j        dk    r!t	          j        | j        | j        |          S | j        j        dvrt          | |          S | 
                                }d|_        |                                 |                    d          r]t          j        r|j        dk    rF|                    |          }|d	k    r|S |dk    rt	          j        |          S t%          d
          |dk    r|                    |j                  j        S |d	k    r|	                    |          S t%          d
          )a  Returns a dense matrix representing the same value.

        Args:
            order ({'C', 'F', None}): Whether to store data in C (row-major)
                order or F (column-major) order. Default is C-order.
            out: Not supported.

        Returns:
            cupy.ndarray: Dense array representing the same matrix.

        .. seealso:: :meth:`scipy.sparse.csr_matrix.toarray`

        r   r)   NC)r   r8   orderfdFDFsparseToDenseFzorder not understood)r-   r*   upperr   rB   rI   r   r8   r   	csr2denser   has_canonical_formatr.   r0   r   r}   r   ascontiguousarrayr   	csc2denser~   )r   r   rX   r*   r!   r&   s         r   rS   zcsr_matrix.toarray  s\    	#"""""}%++--8q==:DJdjNNNN:?&((T5)))IIKK!&	''88 	9 	9,-EAII&&q))A||#-a000 !7888 ||))!#..00#))!,,, !7888r   Fc                     t           r   r   )r   	blocksizer   s      r   tobsrzcsr_matrix.tobsr      !!r   c                     ddl m} |r3| j                                        }| j                                        }n| j        }| j        }|                    | ||          S )zConverts the matrix to COOrdinate format.

        Args:
            copy (bool): If ``False``, it shares data arrays as much as
                possible.

        Returns:
            cupyx.scipy.sparse.coo_matrix: Converted matrix.

        r   r)   )r-   r*   r   r   r   csr2coo)r   r   r*   r   r   s        r   r   zcsr_matrix.tocoo  sj     	#""""" 	#9>>##Dl''))GG9DlGdG444r   c                     ddl m} |                    d          r|j        }n$|                    d          r|j        }nt
           ||           S )a{  Converts the matrix to Compressed Sparse Column format.

        Args:
            copy (bool): If ``False``, it shares data arrays as much as
                possible. Actually this option is ignored because all
                arrays in a matrix cannot be shared in csr to csc conversion.

        Returns:
            cupyx.scipy.sparse.csc_matrix: Converted matrix.

        r   r)   csr2csc
csr2cscEx2)r-   r*   r0   r   r   r1   )r   r   r*   r   s       r   tocsczcsr_matrix.tocsc  sl     	#""""" &&y11 	&&GG((66 	&)GG%%wt}}r   c                 2    |r|                                  S | S )a  Converts the matrix to Compressed Sparse Row format.

        Args:
            copy (bool): If ``False``, the method returns itself.
                Otherwise it makes a copy of the matrix.

        Returns:
            cupyx.scipy.sparse.csr_matrix: Converted matrix.

        r   r   r   s     r   r/   zcsr_matrix.tocsr  s      	99;;Kr   c                 *    |                                  S )zInverts the format.
        )r   )r   s    r   _tocsxzcsr_matrix._tocsx  s     zz||r   c                     t           r   r   r  s     r   todiazcsr_matrix.todia  r   r   c                     t           r   r   r  s     r   todokzcsr_matrix.todok  r   r   c                     t           r   r   r  s     r   tolilzcsr_matrix.tolil  r   r   c                     |t          d          | j        d         | j        d         f}t          j        | j        | j        | j        f||          }| j        |_        |S )aT  Returns a transpose matrix.

        Args:
            axes: This option is not supported.
            copy (bool): If ``True``, a returned matrix shares no data.
                Otherwise, it shared data arrays as much as possible.

        Returns:
            cupyx.scipy.sparse.csc_matrix: `self` with the dimensions reversed.

        NzoSparse matrices do not support an 'axes' parameter because swapping dimensions is the only logical permutation.r9   r   )r   r   )r   r   r   
csc_matrixr   r   r   r   )r   axesr   r   transs        r   	transposezcsr_matrix.transpose  sz     GH H H 
1tz!},Ydk2%dL L L%)%>"r   c                 R    |                      t          ||dz             d          S )zReturns a copy of row i of the matrix, as a (1 x n)
        CSR matrix (row vector).

        Args:
            i (integer): Row

        Returns:
            cupyx.scipy.sparse.csr_matrix: Sparse matrix with single row
        r9   Tr   )_major_sliceslicer   r   s     r   getrowzcsr_matrix.getrow  )       q!a%t <<<r   c                 R    |                      t          ||dz             d          S )zReturns a copy of column i of the matrix, as a (m x 1)
        CSR matrix (column vector).

        Args:
            i (integer): Column

        Returns:
            cupyx.scipy.sparse.csr_matrix: Sparse matrix with single column
        r9   Tr   )_minor_slicer  r  s     r   getcolzcsr_matrix.getcol*  r  r   c                 x    t          ||dz             }|                     |                              |          S Nr9   )r  r  _minor_index_fancyr   r   r   s      r   _get_intXarrayzcsr_matrix._get_intXarray6  s7    Cq!!  %%88===r   c                 |    t          ||dz             }|                     |                              |d          S )Nr9   Tr   )r  r  r  r  s      r   _get_intXslicezcsr_matrix._get_intXslice:  s<    Cq!!  %%223T2BBBr   c                     t          ||dz             }|j        dv }|                     |                              ||          S )Nr9   r9   Nr   )r  stepr  r  )r   r   r   r   s       r   _get_sliceXintzcsr_matrix._get_sliceXint>  sH    Cq!!x9$  %%223T2BBBr   c                 R    |                      |                              |          S r   )r  r  r  s      r   _get_sliceXarrayzcsr_matrix._get_sliceXarrayC  s$      %%88===r   c                 x    t          ||dz             }|                     |                              |          S r  )r  _major_index_fancyr  r  s      r   _get_arrayXintzcsr_matrix._get_arrayXintF  s7    Cq!!&&s++88===r   c                    |j         dvr[|                    | j        d                   \  }}}t          j        |||| j        j                  }|                     ||          S |                     |                              |          S )Nr!  r9   )	r"  r   r   rB   rK   r8   _get_arrayXarrayr'  r  )r   r   r   startstopr"  r   s          r   _get_arrayXslicezcsr_matrix._get_arrayXsliceJ  s    89$$ #DJqM : :E4;udD$,2DEED((d333&&s++88===r   r   )r   )NN)NF)F)/__name__
__module____qualname____doc__formatr   r#   r'   r5   rY   r^   rb   re   rh   rk   rn   r   r   r   r   r   r   r   r   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   r   r      s        : F8 8 8 8&+ + +  1 1 1!" !" !"F< < << < << < << < << < << < <O" O" O"b" " "" " "! ! !F  
	 	 	 	& & &" " ">6 6 66 6 6! ! !$   0+ + +,9 ,9 ,9 ,9\" " " "5 5 5 5,   0      
" " " "" " " "" " " "   .
= 
= 
=
= 
= 
=> > >C C CC C C
> > >> > >> > > > >r   r   c                 ,    t          | t                    S )zChecks if a given matrix is of CSR format.

    Returns:
        bool: Returns if ``x`` is :class:`cupyx.scipy.sparse.csr_matrix`.

    )
isinstancer   r   s    r   rO   rO   R  s     a$$$r   c                     |rR| \  }}|\  }}||k    s|dk    s|dk    st          d          ||k    s|dk    s|dk    st          d          d S d S d S | |k    rt          d          d S )Nr9   zinconsistent shape)r   )a_shapeb_shaper   a_ma_nb_mb_ns          r   r   r   \  s     	3SSs

cQhh#((1222s

cQhh#((1222 
hh(( g1222 r   c                     | j         |z  }| j                                        }| j                                        }t	          |||f| j                  S )Nr   )r   r   r   r   r   r   )spar   r   r   s        r   r   r   i  sK    7Q;DjooGY^^FtWf-RX>>>>r   c                     t          | j        |j                   | j        \  }}|j        \  }}t          ||          t          ||          }}| j        ||z  z  ||z  z  }t	          j        | j        |j                  }	t          j        ||	          }
t          j        || j	        j                  }||k    rW||k    r&t          j
        d|dz   || j        j                  }nOt          j
        d|dz   | j        | j        j                  }n$| j                                        }||k    r||z  } t                      | j        | j        | j	        |||||||||
|           t          |
||f||f          S )Nr7   r   r9   r   )r   r   r   r   rE   r   r8   rB   r   r   rK   r   r   cupy_multiply_by_denser   r   )r=  dnsp_msp_ndn_mdn_nr"   nr   r8   r   r   r   s                r   r   r   p  s}    28444JD$JD$tT??CdOOqA
&AI
!t)
,C"(33E:c'''DjBJ$4555G4xxt88[CE1BIODDDFF[CE26IIIFF!!t88aKF RWbiT4tVQ4J J J tWf-aV<<<<r   al  
__device__ inline int get_row_id(int i, int min, int max, const int *indptr) {
    int row = (min + max) / 2;
    while (min < max) {
        if (i < indptr[row]) {
            max = row - 1;
        } else if (i >= indptr[row + 1]) {
            min = row + 1;
        } else {
            break;
        }
        row = (min + max) / 2;
    }
    return row;
}
a  
__device__ inline int find_index_holding_col_in_row(
        int row, int col, const int *indptr, const int *indices) {
    int j_min = indptr[row];
    int j_max = indptr[row+1] - 1;
    while (j_min <= j_max) {
        int j = (j_min + j_max) / 2;
        int j_col = indices[j];
        if (j_col == col) {
            return j;
        } else if (j_col < col) {
            j_min = j + 1;
        } else {
            j_max = j - 1;
        }
    }
    return -1;
}
)for_each_devicec                  >    t          j        ddddt                    S )Nz
        raw S SP_DATA, raw I SP_INDPTR, raw I SP_INDICES,
        int32 SP_M, int32 SP_N,
        raw D DN_DATA, int32 DN_M, int32 DN_N,
        raw I OUT_INDPTR, int32 OUT_M, int32 OUT_N
        zO OUT_DATA, I OUT_INDICESa  
        int i_out = i;
        int m_out = get_row_id(i_out, 0, OUT_M - 1, &(OUT_INDPTR[0]));
        int i_sp = i_out;
        if (OUT_M > SP_M && SP_M == 1) {
            i_sp -= OUT_INDPTR[m_out];
        }
        if (OUT_N > SP_N && SP_N == 1) {
            i_sp /= OUT_N;
        }
        int n_out = SP_INDICES[i_sp];
        if (OUT_N > SP_N && SP_N == 1) {
            n_out = i_out - OUT_INDPTR[m_out];
        }
        int m_dn = m_out;
        if (OUT_M > DN_M && DN_M == 1) {
            m_dn = 0;
        }
        int n_dn = n_out;
        if (OUT_N > DN_N && DN_N == 1) {
            n_dn = 0;
        }
        OUT_DATA = (O)(SP_DATA[i_sp] * DN_DATA[n_dn + (DN_N * m_dn)]);
        OUT_INDICES = n_out;
        (cupyx_scipy_sparse_csr_multiply_by_densepreamblerB   ElementwiseKernel_GET_ROW_ID_r%   r   r   r@  r@    s4    !	 	$	2 	3E# # # #r   c                  0    t          j        dddd          S )Nz*T data, I row, I col, I width, raw T otherzT resz7
        res = data / other[row * width + col]
        #cupyx_scipy_sparse_coo_divide_denserB   rM  r%   r   r   r   r     s'    !4	 	.  r   c                 T   t          | j        |j                   | j        \  }}|j        \  }}t          ||          t          ||          }}| j        ||z  z  ||z  z  }|j        ||z  z  ||z  z  }	||	k    rt	          ||           S |}
t          j        | j        |j                  }t          j	        |
|          }t          j	        |
| j
        j                  }||k    rW||k    r&t          j        d|
dz   || j        j                  }nOt          j        d|
dz   | j        | j        j                  }n$| j                                        }||k    r||z  }t          j        |
dz   | j
        j                  }t          j        |dz   | j        j                  } t                      | j        | j        | j
        |||j        |j        |j
        |||||||||           t          j        || j        j                  }t          j        || j        j                  }t%          |d                   }t          j	        ||          }t          j	        || j
        j                  } t'                      |||||           t)          |||f||f          S )Nr7   r   r9   r   )r   r   r   r   r   rE   r   r8   rB   r   r   rK   r   r   rI   cupy_multiply_by_csr_step1r   cumsumintcupy_multiply_by_csr_step2r   )r>  r   r8  r9  r:  r;  r"   rF  a_nnzb_nnzc_nnzr8   c_data	c_indicesc_indptrr   nnz_each_rowd_indptrd_nnzd_data	d_indicess                        r   r   r     s    !'222wHCwHCsC==#c3--qAEQ#X!s(+EEQ#X!s(+Eu}}q!$$$E11EZU+++F
5	888I3wws77{1eAgqGGGHH{1eAgquAHNKKKHH8==??s77MHJuQwaio666E:ac888L !  	!)S#	!)S#!Q	5,@ @ @
 KQX^444E{<qx~>>>HEZU+++F
5	888I !  E69MMMvy(3Aq6BBBBr   c                  N    t          j        ddddt          t          z             S )Nz
        raw A A_DATA, raw I A_INDPTR, raw I A_INDICES, int32 A_M, int32 A_N,
        raw B B_DATA, raw I B_INDPTR, raw I B_INDICES, int32 B_M, int32 B_N,
        raw I C_INDPTR, int32 C_M, int32 C_N
        z6C C_DATA, I C_INDICES, raw I FLAGS, raw I NNZ_EACH_ROWaW  
        int i_c = i;
        int m_c = get_row_id(i_c, 0, C_M - 1, &(C_INDPTR[0]));

        int i_a = i;
        if (C_M > A_M && A_M == 1) {
            i_a -= C_INDPTR[m_c];
        }
        if (C_N > A_N && A_N == 1) {
            i_a /= C_N;
        }
        int n_c = A_INDICES[i_a];
        if (C_N > A_N && A_N == 1) {
            n_c = i % C_N;
        }
        int m_b = m_c;
        if (C_M > B_M && B_M == 1) {
            m_b = 0;
        }
        int n_b = n_c;
        if (C_N > B_N && B_N == 1) {
            n_b = 0;
        }
        int i_b = find_index_holding_col_in_row(m_b, n_b,
            &(B_INDPTR[0]), &(B_INDICES[0]));
        if (i_b >= 0) {
            atomicAdd(&(NNZ_EACH_ROW[m_c+1]), 1);
            FLAGS[i+1] = 1;
            C_DATA = (C)(A_DATA[i_a] * B_DATA[i_b]);
            C_INDICES = n_c;
        }
        ,cupyx_scipy_sparse_csr_multiply_by_csr_step1rJ  )rB   rM  rN  _FIND_INDEX_HOLDING_COL_IN_ROW_r%   r   r   rT  rT    s<    !	
 	A	@ 	7 ??Q) ) ) )r   c                  0    t          j        dddd          S )Nz"T C_DATA, I C_INDICES, raw I FLAGSzraw D D_DATA, raw I D_INDICESz
        int j = FLAGS[i];
        if (j < FLAGS[i+1]) {
            D_DATA[j] = (D)(C_DATA);
            D_INDICES[j] = C_INDICES;
        }
        ,cupyx_scipy_sparse_csr_multiply_by_csr_step2rQ  r%   r   r   rW  rW  =  s'    !,'	 	7  r   zH
__device__ inline O binopt(T in1, T in2) {
    return max(in1, in2);
}
zH
__device__ inline O binopt(T in1, T in2) {
    return min(in1, in2);
}
zG
__device__ inline O binopt(T in1, T in2) {
    return (in1 == in2);
}
zG
__device__ inline O binopt(T in1, T in2) {
    return (in1 != in2);
}
zF
__device__ inline O binopt(T in1, T in2) {
    return (in1 < in2);
}
zF
__device__ inline O binopt(T in1, T in2) {
    return (in1 > in2);
}
zG
__device__ inline O binopt(T in1, T in2) {
    return (in1 <= in2);
}
zG
__device__ inline O binopt(T in1, T in2) {
    return (in1 >= in2);
}
c                    t          | j        |j                   | j        \  }}|j        \  }}t          ||          t          ||          }}| j        ||z  z  ||z  z  }	|j        ||z  z  ||z  z  }
t	          j        |	dz   | j        j                  }t	          j        |
dz   |j        j                  }t	          j        |	t          j	                  }t	          j        |
t          j	                  }t	          j        |dz   | j
        j                  }t          j        | j        |j                  }| j                            |d          }|j                            |d          }t          }|dk    r|t          z  }|}n|dk    r|t           z  }|}n|dk    r|t"          z  }t          j        }n|dk    r|t&          z  }t          j        }n|d	k    r|t(          z  }t          j        }ny|d
k    r|t*          z  }t          j        }n\|dk    r|t,          z  }t          j        }n?|dk    r|t.          z  }t          j        }n"t1          d                    |                    t	          j        |	|          }t	          j        |
|          }t	          j        |	| j        j                  }t	          j        |
|j        j                  }|	|
z   } t7          ||          ||| j
        | j        |||| j        |	|j
        |j        ||||j        |
||||||||||           t	          j        ||j                  }t	          j        ||j                  }t	          j        ||j                  }t;          |d                   }t	          j        || j        j                  }t	          j        ||          } t=          |          |||||	|||||
|||           t?          |||f||f          S )Nr9   r7   Fr   r   r   r>   r:   r<   r=   r?   r@   zinvalid op_name: {}rJ  )r   rS  r   ) r   r   r   r   rB   rI   r   r8   rE   int8r   r   r   r   rN  _BINOPT_MAX__BINOPT_MIN__BINOPT_EQ_rH   _BINOPT_NE__BINOPT_LT__BINOPT_GT__BINOPT_LE__BINOPT_GE_r   r2  r   cupy_binopt_csr_step1rU  rV  cupy_binopt_csr_step2r   )r>  r   rU   r8  r9  r:  r;  r"   rF  rX  rY  a_infob_infoa_validb_validr]  in_dtypea_datab_datafuncs	out_dtype
a_tmp_data
b_tmp_dataa_tmp_indicesb_tmp_indices_sizerZ  r\  r[  s                                r   rL   rL   w  s    !'222wHCwHCsC==#c3--qAEQ#X!s(+EEQ#X!s(+EZ	999FZ	999Fjej111Gjej111Gz!a%qx~666H"17AG44HV]]8%]00FV]]8%]00FE+			K					F		K			F		K			F		K			F		K			F		K			F		K		.55g>>???E333JE333JJuAIO<<<MJuAIO<<<MEME2'E222	1	!)VS#que	!)VS#que

u    [v|444F[v|444F{88>:::HE
5	888IZY///F"'""
E
E6' ' ' ' vy(3Aq6BBBBr    c                 D    d| z   dz   }t          j        ddd||          S )Ncupyx_scipy_sparse_csr_binopt_step1a  
        int32 M, int32 N,
        raw I A_INDPTR, raw I A_INDICES, raw T A_DATA,
        int32 A_M, int32 A_N, int32 A_NNZ_ACT, int32 A_NNZ,
        raw I B_INDPTR, raw I B_INDICES, raw T B_DATA,
        int32 B_M, int32 B_N, int32 B_NNZ_ACT, int32 B_NNZ
        z
        raw I A_INFO, raw B A_VALID, raw I A_TMP_INDICES, raw O A_TMP_DATA,
        raw I B_INFO, raw B B_VALID, raw I B_TMP_INDICES, raw O B_TMP_DATA,
        raw I C_INFO
        ab  
        if (i >= A_NNZ + B_NNZ) return;

        const int *MY_INDPTR, *MY_INDICES;  int *MY_INFO;  const T *MY_DATA;
        const int *OP_INDPTR, *OP_INDICES;  int *OP_INFO;  const T *OP_DATA;
        int MY_M, MY_N, MY_NNZ_ACT, MY_NNZ;
        int OP_M, OP_N, OP_NNZ_ACT, OP_NNZ;
        signed char *MY_VALID;  I *MY_TMP_INDICES;  O *MY_TMP_DATA;

        int my_j;
        if (i < A_NNZ) {
            // in charge of one of non-zero element of sparse matrix A
            my_j = i;
            MY_INDPTR  = &(A_INDPTR[0]);   OP_INDPTR  = &(B_INDPTR[0]);
            MY_INDICES = &(A_INDICES[0]);  OP_INDICES = &(B_INDICES[0]);
            MY_INFO    = &(A_INFO[0]);     OP_INFO    = &(B_INFO[0]);
            MY_DATA    = &(A_DATA[0]);     OP_DATA    = &(B_DATA[0]);
            MY_M       = A_M;              OP_M       = B_M;
            MY_N       = A_N;              OP_N       = B_N;
            MY_NNZ_ACT = A_NNZ_ACT;        OP_NNZ_ACT = B_NNZ_ACT;
            MY_NNZ     = A_NNZ;            OP_NNZ     = B_NNZ;
            MY_VALID   = &(A_VALID[0]);
            MY_TMP_DATA= &(A_TMP_DATA[0]);
            MY_TMP_INDICES = &(A_TMP_INDICES[0]);
        } else {
            // in charge of one of non-zero element of sparse matrix B
            my_j = i - A_NNZ;
            MY_INDPTR  = &(B_INDPTR[0]);   OP_INDPTR  = &(A_INDPTR[0]);
            MY_INDICES = &(B_INDICES[0]);  OP_INDICES = &(A_INDICES[0]);
            MY_INFO    = &(B_INFO[0]);     OP_INFO    = &(A_INFO[0]);
            MY_DATA    = &(B_DATA[0]);     OP_DATA    = &(A_DATA[0]);
            MY_M       = B_M;              OP_M       = A_M;
            MY_N       = B_N;              OP_N       = A_N;
            MY_NNZ_ACT = B_NNZ_ACT;        OP_NNZ_ACT = A_NNZ_ACT;
            MY_NNZ     = B_NNZ;            OP_NNZ     = A_NNZ;
            MY_VALID   = &(B_VALID[0]);
            MY_TMP_DATA= &(B_TMP_DATA[0]);
            MY_TMP_INDICES = &(B_TMP_INDICES[0]);
        }
        int _min, _max, _mid;

        // get column location
        int my_col;
        int my_j_act = my_j;
        if (MY_M == 1 && MY_M < M) {
            if (MY_N == 1 && MY_N < N) my_j_act = 0;
            else                       my_j_act = my_j % MY_NNZ_ACT;
        } else {
            if (MY_N == 1 && MY_N < N) my_j_act = my_j / N;
        }
        my_col = MY_INDICES[my_j_act];
        if (MY_N == 1 && MY_N < N) {
            my_col = my_j % N;
        }

        // get row location
        int my_row = get_row_id(my_j_act, 0, MY_M - 1, &(MY_INDPTR[0]));
        if (MY_M == 1 && MY_M < M) {
            if (MY_N == 1 && MY_N < N) my_row = my_j / N;
            else                       my_row = my_j / MY_NNZ_ACT;
        }

        int op_row = my_row;
        int op_row_act = op_row;
        if (OP_M == 1 && OP_M < M) {
            op_row_act = 0;
        }

        int op_col = 0;
        _min = OP_INDPTR[op_row_act];
        _max = OP_INDPTR[op_row_act + 1] - 1;
        int op_j_act = _min;
        bool op_nz = false;
        if (_min <= _max) {
            if (OP_N == 1 && OP_N < N) {
                op_col = my_col;
                op_nz = true;
            }
            else {
                _mid = (_min + _max) / 2;
                op_col = OP_INDICES[_mid];
                while (_min < _max) {
                    if (op_col < my_col) {
                        _min = _mid + 1;
                    } else if (op_col > my_col) {
                        _max = _mid;
                    } else {
                        break;
                    }
                    _mid = (_min + _max) / 2;
                    op_col = OP_INDICES[_mid];
                }
                op_j_act = _mid;
                if (op_col == my_col) {
                    op_nz = true;
                } else if (op_col < my_col) {
                    op_col = N;
                    op_j_act += 1;
                }
            }
        }

        int op_j = op_j_act;
        if (OP_M == 1 && OP_M < M) {
            if (OP_N == 1 && OP_N < N) {
                op_j = (op_col + N * op_row) * OP_NNZ_ACT;
            } else {
                op_j = op_j_act + OP_NNZ_ACT * op_row;
            }
        } else {
            if (OP_N == 1 && OP_N < N) {
                op_j = op_col + N * op_j_act;
            }
        }

        if (i < A_NNZ || !op_nz) {
            T my_data = MY_DATA[my_j_act];
            T op_data = 0;
            if (op_nz) op_data = OP_DATA[op_j_act];
            O out;
            if (i < A_NNZ) out = binopt(my_data, op_data);
            else           out = binopt(op_data, my_data);
            if (out != static_cast<O>(0)) {
                MY_VALID[my_j] = 1;
                MY_TMP_DATA[my_j] = out;
                MY_TMP_INDICES[my_j] = my_col;
                atomicAdd( &(C_INFO[my_row + 1]), 1 );
                atomicAdd( &(MY_INFO[my_j + 1]), 1 );
                atomicAdd( &(OP_INFO[op_j]), 1 );
            }
        }
        rJ  rQ  )rU   rK  names      r   rr  rr    sH    +g5?D!		
C	H 	xcR R R Rr   c                 @    d| z   dz   }t          j        ddd|          S )Ncupyx_scipy_sparse_csr_binoptstep2z
        raw I A_INFO, raw B A_VALID, raw I A_TMP_INDICES, raw O A_TMP_DATA,
        int32 A_NNZ,
        raw I B_INFO, raw B B_VALID, raw I B_TMP_INDICES, raw O B_TMP_DATA,
        int32 B_NNZ
        zraw I C_INDICES, raw O C_DATAa  
        if (i < A_NNZ) {
            int j = i;
            if (A_VALID[j]) {
                C_INDICES[A_INFO[j]] = A_TMP_INDICES[j];
                C_DATA[A_INFO[j]]    = A_TMP_DATA[j];
            }
        } else if (i < A_NNZ + B_NNZ) {
            int j = i - A_NNZ;
            if (B_VALID[j]) {
                C_INDICES[B_INFO[j]] = B_TMP_INDICES[j];
                C_DATA[B_INFO[j]]    = B_TMP_DATA[j];
            }
        }
        rQ  )rU   r  s     r   rs  rs  S  s:    *W4w>D!	 	(	 	/  r   c           	          t          j        | j        | j        |          }| j        \  }}t	          | j                  } |||| j        | j        | j        |dk    |           |S )N)r8   r   r   )rB   rI   r   r8   _cupy_csr2denser   r   r   )r>  r   rX   r"   rF  kerns         r   r   r   q  sa    
*QWAG5
9
9
9C7DAq17##DDAqxAFUc\C@@@Jr   c                 Z    | dk    rd}nd}t          j        ddd|z   dt                    S )	N?zif (DATA) OUT[index] = true;zatomicAdd(&OUT[index], DATA);z?int32 M, int32 N, raw I INDPTR, I INDICES, T DATA, bool C_ORDERz	raw T OUTz
        int row = get_row_id(i, 0, M - 1, &(INDPTR[0]));
        int col = INDICES;
        int index = C_ORDER ? col + N * row : row + M * col;
        cupyx_scipy_sparse_csr2denserJ  rL  )r8   rT   s     r   r  r  y  sN    ||+,!I	 		
 	'
 
 
 
r   c                 4   ddl m} | j        j        dv rA|                    d          r|                    | d          S |                    |           S | j        \  }}t          j	        |           } t          j
        |dz   t          j                  }t          j
        ||z  dz   t          j                  } t                      ||| ||           t          j        |t          j                  }t          j        |t          j                  }t          |d	                   }t          j        |t          j                  }t          j        || j                  } t#                      ||| |||           t%          |||f||f
          S )Nr   r)   r   denseToSparser   )r2  r9   r7   rS  r   )r-   r*   r8   r   r0   r  r    r   rB   r   rI   rE   rJ   cupy_dense2csr_step1rU  rV  r   cupy_dense2csr_step2r   )	r>  r*   r"   rF  r   infor   r   r   s	            r   r    r      s~   w|v&&77 	)))!E):::%%a(((7DAqq!!AZAU[111F:a!eaiu{333D1aFD111[u{333F;t5;///D
fRj//CjEK000G:c)))D1aD'4888tWf-aV<<<<r   c                  0    t          j        dddd          S )Nzint32 M, int32 N, T Azraw I INDPTR, raw I INFOz
        int row = i / N;
        int col = i % N;
        if (A != static_cast<T>(0)) {
            atomicAdd( &(INDPTR[row + 1]), 1 );
            INFO[i + 1] = 1;
        }
        "cupyx_scipy_sparse_dense2csr_step1rQ  r%   r   r   r  r    s'    !"	 	-. . .r   c                  0    t          j        dddd          S )Nz!int32 M, int32 N, T A, raw I INFOzraw I INDICES, raw T DATAz
        int row = i / N;
        int col = i % N;
        if (A != static_cast<T>(0)) {
            int idx = INFO[i];
            INDICES[idx] = col;
            DATA[idx] = A;
        }
        "cupyx_scipy_sparse_dense2csr_step2rQ  r%   r   r   r  r    s'    !+#	 	-. . .r   c                  >    t          j        ddddt                    S )NzHint32 k, int32 rows, int32 cols, raw T data, raw I indptr, raw I indiceszT yal  
        int row = i;
        int col = i;
        if (k < 0) row -= k;
        if (k > 0) col += k;
        if (row >= rows || col >= cols) return;
        int j = find_index_holding_col_in_row(row, col,
            &(indptr[0]), &(indices[0]));
        if (j >= 0) {
            y = data[j];
        } else {
            y = static_cast<T>(0);
        }
        cupyx_scipy_sparse_csr_diagonalrJ  )rB   rM  re  r%   r   r   r   r     s1    !	2	 	*0'   r   )T)r  )2r[   rP   rE   scipy.sparser   r   ImportErrorrB   
cupy._corer   	cupy.cudar   r   cupyx.scipy.sparser   r   r   r	   r
   _compressed_sparse_matrixr   rO   r   r   r   rN  re  memoizer@  r   r   rT  rW  rj  rk  rl  rm  rn  ro  rp  rq  rL   rr  rs  r   r  r    r  r  r   r%   r   r   <module>r     s          # # # # # #             $ $ $ $ $ $ * * * * * * # # # # # # 6 6 6 6 6 6 $ $ $ $ $ $x> x> x> x> x>6 x> x> x>v% % %
3 
3 
3 
3? ? ?= = =4"# * D))$ $ *)$N D))  *)(C (C (CV D))* * *)*Z D))  *)






AC AC ACH D))T T T *)Tn D))  *):   D))  *)&= = =, D)). . *). D)). . *).  D))  *)  s    