Repository: cupy/cupy
Issue: cupyx.scipy.sparse.csr_matrix of dtype bool behaves different from scipy.sparse.csr_matrix

### Description

Both `cupyx.scipy.sparse.csr_matrix` and `scipy.sparse.csr_matrix` support non-canonical data, i.e. unsorted column indices with duplicate entries. For dtype `bool` their behavior with duplicates is different: `scipy.sparse.csr_matrix` adds the values, while `cupyx.scipy.sparse.csr_matrix` seems to take the first entry for a specific coordinate. See example reproducer!

### To Reproduce

```py
import numpy as np
import scipy.sparse as sp
import cupy as cp
import cupyx.scipy.sparse as csp

data = np.array((True, True, True, False, True))
indices = np.array((1, 0, 0, 1, 1), dtype=int)
indptr = np.array((0, 2, 5), dtype=int)

csr = sp.csr_matrix((data, indices, indptr))
print(csr.todense())

# [[ True  True]
#  [ True  True]]

c_csr = csp.csr_matrix((cp.array(data), cp.array(indices), cp.array(indptr)))
print(c_csr.todense())

# [[ True  True]
#  [ True False]]
```


### Installation

Source (`pip install cupy`)

### Environment

```
(cupytest311) C:\\Users\\weber\\Documents\\src\\sparseconverter>python -c "import cupy; cupy.show_config()"
OS                           : Windows-10-10.0.22621-SP0
Python Version               : 3.11.4
CuPy Version                 : 13.0.0a1
CuPy Platform                : NVIDIA CUDA
NumPy Version                : 1.24.4
SciPy Version                : 1.11.1
Cython Build Version         : 0.29.36
Cython Runtime Version       : None
CUDA Root                    : C:\\Program Files\\NVIDIA GPU Computing Toolkit\\CUDA\\v12.1
nvcc PATH                    : C:\\Program Files\\NVIDIA GPU Computing Toolkit\\CUDA\\v12.1\\bin\\nvcc.EXE
CUDA Build Version           : 12010
CUDA Driver Version          : 12010
CUDA Runtime Version         : 12010
cuBLAS Version               : (available)
cuFFT Version                : 11002
cuRAND Version               : 10302
cuSOLVER Version             : (11, 4, 5)
cuSPARSE Version             : (available)
NVRTC Version                : (12, 1)
Thrust Version               : 200001
CUB Build Version            : 200001
Jitify Build Version         : 4a37de0
cuDNN Build Version          : None
cuDNN Version                : None
NCCL Build Version           : None
NCCL Runtime Version         : None
cuTENSOR Version             : None
cuSPARSELt Build Version     : None
Device 0 Name                : NVIDIA T500
Device 0 Compute Capability  : 75
Device 0 PCI Bus ID          : 0000:01:00.0
```

### Additional Information

_No response_

Changed files:
cupyx/scipy/sparse/_csr.py

diff --git a/cupyx/scipy/sparse/_csr.py b/cupyx/scipy/sparse/_csr.py
index 21669490c..102b4954f 100644
--- a/cupyx/scipy/sparse/_csr.py
+++ b/cupyx/scipy/sparse/_csr.py
@@ -1108,25 +1108,26 @@ def cupy_binopt_csr_step2(op_name):
 def csr2dense(a, order):
     out = cupy.zeros(a.shape, dtype=a.dtype, order=order)
     m, n = a.shape
-    cupy_csr2dense()(m, n, a.indptr, a.indices, a.data,
-                     (order == 'C'), out)
+    kern = _cupy_csr2dense(a.dtype)
+    kern(m, n, a.indptr, a.indices, a.data, (order == 'C'), out)
     return out
 
 
 @cupy._util.memoize(for_each_device=True)
-def cupy_csr2dense():
+def _cupy_csr2dense(dtype):
+    if dtype == '?':
+        op = "if (DATA) OUT[index] = true;"
+    else:
+        op = "atomicAdd(&OUT[index], DATA);"
+
     return cupy.ElementwiseKernel(
         'int32 M, int32 N, raw I INDPTR, I INDICES, T DATA, bool C_ORDER',
         'raw T OUT',
         '''
         int row = get_row_id(i, 0, M - 1, &(INDPTR[0]));
         int col = INDICES;
-        if (C_ORDER) {
-            OUT[col + N * row] += DATA;
-        } else {
-            OUT[row + M * col] += DATA;
-        }
-        ''',
+        int index = C_ORDER ? col + N * row : row + M * col;
+        ''' + op,
         'cupyx_scipy_sparse_csr2dense',
         preamble=_GET_ROW_ID_
     )