Skip to content

Commit 0475e11

Browse files
EgorOrachyovEgorOrachyov
committed
[Code] Add matrix row/col extraction methods for cuda && seq backend
1 parent 8605dbf commit 0475e11

17 files changed

Lines changed: 482 additions & 13 deletions

‎cubool/CMakeLists.txt‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -75,6 +75,8 @@ set(CUBOOL_C_API_SOURCES
7575
sources/cuBool_Matrix_Marker.cpp
7676
sources/cuBool_Matrix_ExtractPairs.cpp
7777
sources/cuBool_Matrix_ExtractSubMatrix.cpp
78+
sources/cuBool_Matrix_ExtractRow.cpp
79+
sources/cuBool_Matrix_ExtractCol.cpp
7880
sources/cuBool_Matrix_Duplicate.cpp
7981
sources/cuBool_Matrix_Transpose.cpp
8082
sources/cuBool_Matrix_Nvals.cpp

‎cubool/include/cubool/cubool.h‎

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -363,6 +363,48 @@ CUBOOL_EXPORT CUBOOL_API cuBool_Status cuBool_Matrix_ExtractSubMatrix(
363363
cuBool_Hints hints
364364
);
365365

366+
/**
367+
* Extract specified matrix row as vector.
368+
*
369+
* @note Vector and matrix must have compatible size.
370+
* dim(matrix) = M x N
371+
* dim(vector) = N
372+
*
373+
* @param result Vector handle where to store extracted row
374+
* @param matrix Source matrix
375+
* @param i Index of the matrix row to extract
376+
* @param hints Hints for the operation
377+
*
378+
* @return Error code on this operation
379+
*/
380+
CUBOOL_EXPORT CUBOOL_API cuBool_Status cuBool_Matrix_ExtractRow(
381+
cuBool_Vector result,
382+
cuBool_Matrix matrix,
383+
cuBool_Index i,
384+
cuBool_Hints hints
385+
);
386+
387+
/**
388+
* Extract specified matrix col as vector.
389+
*
390+
* @note Vector and matrix must have compatible size.
391+
* dim(matrix) = M x N
392+
* dim(vector) = M
393+
*
394+
* @param result Vector handle where to store extracted column
395+
* @param matrix Source matrix
396+
* @param j Index of the matrix column to extract
397+
* @param hints Hints for the operation
398+
*
399+
* @return Error code on this operation
400+
*/
401+
CUBOOL_EXPORT CUBOOL_API cuBool_Status cuBool_Matrix_ExtractCol(
402+
cuBool_Vector result,
403+
cuBool_Matrix matrix,
404+
cuBool_Index j,
405+
cuBool_Hints hints
406+
);
407+
366408
/**
367409
* Creates new sparse matrix, duplicates content and stores handle in the provided pointer.
368410
*

‎cubool/sources/backend/vector_base.hpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -40,6 +40,8 @@ namespace cubool {
4040
virtual void build(const index *rows, size_t nvals, bool isSorted, bool noDuplicates) = 0;
4141
virtual void extract(index* rows, size_t &nvals) = 0;
4242
virtual void extractSubVector(const VectorBase &otherBase, index i, index nrows, bool checkTime) = 0;
43+
virtual void extractRow(const class MatrixBase& matrixBase, index i) = 0;
44+
virtual void extractCol(const class MatrixBase& matrixBase, index j) = 0;
4345

4446
virtual void clone(const VectorBase& otherBase) = 0;
4547
virtual void reduce(index &result, bool checkTime) = 0;

‎cubool/sources/core/vector.cpp‎

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -115,6 +115,30 @@ namespace cubool {
115115
mHnd->extractSubVector(*other->mHnd, i, nrows, false);
116116
}
117117

118+
void Vector::extractRow(const class MatrixBase &matrixBase, index i) {
119+
const auto* matrix = dynamic_cast<const Matrix*>(&matrixBase);
120+
121+
CHECK_RAISE_ERROR(matrix != nullptr, InvalidArgument, "Passed matrix does not belong to core matrix class");
122+
CHECK_RAISE_ERROR(i < matrix->getNrows(), InvalidArgument, "Row index must be within matrix bounds");
123+
124+
matrix->commitCache();
125+
this->releaseCache();
126+
127+
mHnd->extractRow(*matrix->mHnd, i);
128+
}
129+
130+
void Vector::extractCol(const class MatrixBase &matrixBase, index j) {
131+
const auto* matrix = dynamic_cast<const Matrix*>(&matrixBase);
132+
133+
CHECK_RAISE_ERROR(matrix != nullptr, InvalidArgument, "Passed matrix does not belong to core matrix class");
134+
CHECK_RAISE_ERROR(j < matrix->getNcols(), InvalidArgument, "Column index must be within matrix bounds");
135+
136+
matrix->commitCache();
137+
this->releaseCache();
138+
139+
mHnd->extractCol(*matrix->mHnd, j);
140+
}
141+
118142
void Vector::clone(const VectorBase &otherBase) {
119143
const auto* other = dynamic_cast<const Vector*>(&otherBase);
120144

‎cubool/sources/core/vector.hpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -43,6 +43,8 @@ namespace cubool {
4343
void build(const index *rows, size_t nvals, bool isSorted, bool noDuplicates) override;
4444
void extract(index *rows, size_t &nvals) override;
4545
void extractSubVector(const VectorBase &otherBase, index i, index nrows, bool checkTime) override;
46+
void extractRow(const class MatrixBase& matrixBase, index i) override;
47+
void extractCol(const class MatrixBase& matrixBase, index j) override;
4648

4749
void clone(const VectorBase &otherBase) override;
4850
void reduce(index &result, bool checkTime) override;

‎cubool/sources/cuBool_Matrix_ExtractCol.cpp‎

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
/**********************************************************************************/
2+
/* MIT License */
3+
/* */
4+
/* Copyright (c) 2020, 2021 JetBrains-Research */
5+
/* */
6+
/* Permission is hereby granted, free of charge, to any person obtaining a copy */
7+
/* of this software and associated documentation files (the "Software"), to deal */
8+
/* in the Software without restriction, including without limitation the rights */
9+
/* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell */
10+
/* copies of the Software, and to permit persons to whom the Software is */
11+
/* furnished to do so, subject to the following conditions: */
12+
/* */
13+
/* The above copyright notice and this permission notice shall be included in all */
14+
/* copies or substantial portions of the Software. */
15+
/* */
16+
/* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR */
17+
/* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, */
18+
/* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE */
19+
/* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER */
20+
/* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, */
21+
/* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE */
22+
/* SOFTWARE. */
23+
/**********************************************************************************/
24+
25+
#include <cuBool_Common.hpp>
26+
27+
cuBool_Status cuBool_Matrix_ExtractCol(
28+
cuBool_Vector result,
29+
cuBool_Matrix matrix,
30+
cuBool_Index j,
31+
cuBool_Hints hints
32+
) {
33+
CUBOOL_BEGIN_BODY
34+
CUBOOL_ARG_NOT_NULL(result)
35+
CUBOOL_ARG_NOT_NULL(matrix)
36+
auto r = (cubool::Vector*) result;
37+
auto m = (cubool::Matrix*) matrix;
38+
r->extractCol(*m, j);
39+
CUBOOL_END_BODY
40+
}

‎cubool/sources/cuBool_Matrix_ExtractRow.cpp‎

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
/**********************************************************************************/
2+
/* MIT License */
3+
/* */
4+
/* Copyright (c) 2020, 2021 JetBrains-Research */
5+
/* */
6+
/* Permission is hereby granted, free of charge, to any person obtaining a copy */
7+
/* of this software and associated documentation files (the "Software"), to deal */
8+
/* in the Software without restriction, including without limitation the rights */
9+
/* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell */
10+
/* copies of the Software, and to permit persons to whom the Software is */
11+
/* furnished to do so, subject to the following conditions: */
12+
/* */
13+
/* The above copyright notice and this permission notice shall be included in all */
14+
/* copies or substantial portions of the Software. */
15+
/* */
16+
/* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR */
17+
/* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, */
18+
/* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE */
19+
/* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER */
20+
/* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, */
21+
/* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE */
22+
/* SOFTWARE. */
23+
/**********************************************************************************/
24+
25+
#include <cuBool_Common.hpp>
26+
27+
cuBool_Status cuBool_Matrix_ExtractRow(
28+
cuBool_Vector result,
29+
cuBool_Matrix matrix,
30+
cuBool_Index i,
31+
cuBool_Hints hints
32+
) {
33+
CUBOOL_BEGIN_BODY
34+
CUBOOL_ARG_NOT_NULL(result)
35+
CUBOOL_ARG_NOT_NULL(matrix)
36+
auto r = (cubool::Vector*) result;
37+
auto m = (cubool::Matrix*) matrix;
38+
r->extractRow(*m, i);
39+
CUBOOL_END_BODY
40+
}

‎cubool/sources/cuda/cuda_vector.cu‎

Lines changed: 92 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -23,14 +23,16 @@
2323
/**********************************************************************************/
2424

2525
#include <cuda/cuda_vector.hpp>
26+
#include <cuda/cuda_matrix.hpp>
27+
#include <cuda/kernels/bin_search.cuh>
2628
#include <core/error.hpp>
2729
#include <utils/data_utils.hpp>
2830
#include <limits>
2931

3032
namespace cubool {
3133

3234
CudaVector::CudaVector(size_t nrows, CudaInstance &instance)
33-
: mVectorImpl(nrows), mInstance(instance) {
35+
: mVectorImpl(nrows), mInstance(instance) {
3436

3537
}
3638

@@ -71,7 +73,7 @@ namespace cubool {
7173
}
7274

7375
void CudaVector::extractSubVector(const VectorBase &otherBase, index i, index nrows, bool checkTime) {
74-
const auto* v = dynamic_cast<const CudaVector*>(&otherBase);
76+
const auto *v = dynamic_cast<const CudaVector *>(&otherBase);
7577

7678
CHECK_RAISE_ERROR(v != nullptr, InvalidArgument, "Passed vector does not belong to cuda vector class");
7779

@@ -92,23 +94,23 @@ namespace cubool {
9294
return;
9395
}
9496

95-
auto& vec = v->mVectorImpl;
97+
auto &vec = v->mVectorImpl;
9698

9799
thrust::device_vector<index, DeviceAlloc<index>> region(2);
98100
thrust::fill_n(region.begin(), 1, std::numeric_limits<index>::max());
99101
thrust::fill_n(region.begin() + 1, 1, 0);
100102

101103
thrust::for_each(thrust::counting_iterator<index>(0), thrust::counting_iterator<index>(vec.m_vals),
102104
[first = region.data(), size = region.data() + 1,
103-
i = i, last = i + nrows, rowIndex = vec.m_rows_index.data()]
104-
__device__ (index id) {
105-
auto rowId = rowIndex[id];
105+
i = i, last = i + nrows, rowIndex = vec.m_rows_index.data()]
106+
__device__(index id) {
107+
auto rowId = rowIndex[id];
106108

107-
if (i <= rowId && rowId < last) {
108-
atomicAdd(size.get(), 1);
109-
atomicMin(first.get(), id);
110-
}
111-
});
109+
if (i <= rowId && rowId < last) {
110+
atomicAdd(size.get(), 1);
111+
atomicMin(first.get(), id);
112+
}
113+
});
112114

113115
index resultSize = region.back();
114116

@@ -122,14 +124,91 @@ namespace cubool {
122124
index firstToCopy = region.front();
123125

124126
VectorImplType::container_type result(resultSize);
125-
thrust::copy(vec.m_rows_index.begin() + firstToCopy, vec.m_rows_index.begin() + firstToCopy + resultSize, result.begin());
127+
thrust::copy(vec.m_rows_index.begin() + firstToCopy, vec.m_rows_index.begin() + firstToCopy + resultSize,
128+
result.begin());
126129

127130
// Update this impl data
128131
mVectorImpl = std::move(VectorImplType(std::move(result), nrows, resultSize));
129132
}
130133

134+
void CudaVector::extractRow(const class MatrixBase &matrixBase, index i) {
135+
auto matrix = dynamic_cast<const CudaMatrix *>(&matrixBase);
136+
137+
CHECK_RAISE_ERROR(matrix != nullptr, InvalidArgument, "Provided matrix does not belongs to cuda matrix class");
138+
139+
assert(getNrows() == matrix->getNcols());
140+
assert(i <= matrix->getNrows());
141+
142+
auto &m = matrix->mMatrixImpl;
143+
144+
index beginOffset = m.m_row_index[i];
145+
index endOffset = m.m_row_index[i + 1];
146+
147+
auto size = endOffset - beginOffset;
148+
auto being = m.m_col_index.begin() + beginOffset;
149+
auto end = m.m_col_index.begin() + endOffset;
150+
151+
VectorImplType::container_type result(size);
152+
thrust::copy(being, end, result.begin());
153+
154+
mVectorImpl = std::move(VectorImplType(std::move(result), m.m_cols, size));
155+
}
156+
157+
void CudaVector::extractCol(const class MatrixBase &matrixBase, index j) {
158+
auto matrix = dynamic_cast<const CudaMatrix *>(&matrixBase);
159+
160+
CHECK_RAISE_ERROR(matrix != nullptr, InvalidArgument, "Provided matrix does not belongs to cuda matrix class");
161+
162+
assert(getNrows() == matrix->getNrows());
163+
assert(j <= matrix->getNcols());
164+
165+
auto &m = matrix->mMatrixImpl;
166+
167+
VectorImplType::container_type nnz(1);
168+
thrust::fill(nnz.begin(), nnz.end(), (index) 0);
169+
170+
thrust::for_each(thrust::counting_iterator<index>(0), thrust::counting_iterator<index>(m.m_rows),
171+
[rowOffset = m.m_row_index.data(), colIndex = m.m_col_index.data(),
172+
j, nnz = nnz.data()]__device__(index i) {
173+
auto size = rowOffset[i + 1] - rowOffset[i];
174+
auto begin = colIndex + rowOffset[i];
175+
auto end = colIndex + rowOffset[i + 1];
176+
177+
auto r = kernels::find<index>(begin, end, j);
178+
179+
if (r != end && *r == j)
180+
atomicAdd(nnz.get(), 1);
181+
}
182+
);
183+
184+
index size = nnz.back();
185+
VectorImplType::container_type result(size);
186+
187+
thrust::fill(nnz.begin(), nnz.end(), (index) 0);
188+
189+
thrust::for_each(thrust::counting_iterator<index>(0), thrust::counting_iterator<index>(m.m_rows),
190+
[rowOffset = m.m_row_index.data(), colIndex = m.m_col_index.data(),
191+
j, nnz = nnz.data(), result = result.data()]__device__(index i) {
192+
auto size = rowOffset[i + 1] - rowOffset[i];
193+
auto begin = colIndex + rowOffset[i];
194+
auto end = colIndex + rowOffset[i + 1];
195+
196+
auto r = kernels::find<index>(begin, end, j);
197+
198+
if (r != end && *r == j) {
199+
auto order = atomicAdd(nnz.get(), 1);
200+
result[order] = i;
201+
}
202+
}
203+
);
204+
205+
thrust::sort(result.begin(), result.end());
206+
207+
mVectorImpl = std::move(VectorImplType(std::move(result), m.m_rows, size));
208+
}
209+
131210
void CudaVector::clone(const VectorBase &otherBase) {
132-
auto other = dynamic_cast<const CudaVector*>(&otherBase);
211+
auto other = dynamic_cast<const CudaVector *>(&otherBase);
133212

134213
CHECK_RAISE_ERROR(other != nullptr, InvalidArgument, "Passed vector does not belong to vector class");
135214
CHECK_RAISE_ERROR(other != this, InvalidArgument, "Vectors must differ");

‎cubool/sources/cuda/cuda_vector.hpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -45,6 +45,8 @@ namespace cubool {
4545
void build(const index *rows, size_t nvals, bool isSorted, bool noDuplicates) override;
4646
void extract(index *rows, size_t &nvals) override;
4747
void extractSubVector(const VectorBase &otherBase, index i, index nrows, bool checkTime) override;
48+
void extractRow(const class MatrixBase& matrixBase, index i) override;
49+
void extractCol(const class MatrixBase& matrixBase, index j) override;
4850

4951
void clone(const VectorBase &otherBase) override;
5052
void reduce(index &result, bool checkTime) override;

‎cubool/sources/cuda/kernels/bin_search.cuh‎

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -54,6 +54,31 @@ namespace cubool {
5454
return range;
5555
}
5656

57+
template<typename IndexType, typename Iter>
58+
__host__ __device__ Iter find(Iter begin, Iter end, IndexType v) {
59+
auto s = thrust::distance(begin, end);
60+
61+
if (s == 0)
62+
return end;
63+
64+
auto l = 0;
65+
auto r = s - 1;
66+
67+
while (l <= r) {
68+
auto k = (l + r) / 2;
69+
auto i = begin + k;
70+
71+
if (*i < v)
72+
l = k + 1;
73+
else if (*i > v)
74+
r = k - 1;
75+
else
76+
return i;
77+
}
78+
79+
return end;
80+
}
81+
5782
}
5883
}
5984

0 commit comments

Comments
 (0)