Skip to content

Commit da913de

Browse files
EgorOrachyovEgorOrachyov
committed
[Code] Add cuda mt-vt reduce impl && expose mt-vt reduce into C API && add respective tests
1 parent 199f342 commit da913de

17 files changed

Lines changed: 429 additions & 40 deletions

‎cubool/CMakeLists.txt‎

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,7 @@ if (CUBOOL_RELEASE)
2525
list(APPEND DEFINES_LIST CUBOOL_RELEASE)
2626
endif()
2727

28-
# Platform checks
28+
# Platform checks
2929
if(APPLE)
3030
list(APPEND DEFINES_LIST CUBOOL_PLATFORM_MACOS)
3131
set(TARGET_FILE_NAME "lib${TARGET_NAME}.dylib")
@@ -82,6 +82,7 @@ set(CUBOOL_C_API_SOURCES
8282
sources/cuBool_Matrix_Ncols.cpp
8383
sources/cuBool_Matrix_Free.cpp
8484
sources/cuBool_Matrix_Reduce.cpp
85+
sources/cuBool_Matrix_Reduce2.cpp
8586
sources/cuBool_Matrix_EWiseAdd.cpp
8687
sources/cuBool_Vector_New.cpp
8788
sources/cuBool_Vector_Build.cpp
@@ -127,12 +128,15 @@ if (CUBOOL_WITH_CUDA)
127128
sources/cuda/cuda_matrix_extract_sub_matrix.cu
128129
sources/cuda/cuda_vector.hpp
129130
sources/cuda/cuda_vector.cu
131+
sources/cuda/cuda_vector_ewiseadd.cu
132+
sources/cuda/cuda_vector_reduce.cu
130133
sources/cuda/details/meta.hpp
131134
sources/cuda/details/sp_vector.hpp
132135
sources/cuda/details/host_allocator.hpp
133136
sources/cuda/details/device_allocator.cuh
134137
sources/cuda/kernels/slow_sort.cuh
135138
sources/cuda/kernels/bin_search.cuh
139+
sources/cuda/kernels/spewiseadd.cuh
136140
sources/cuda/kernels/sptranspose.cuh
137141
sources/cuda/kernels/sptranspose2.cuh
138142
sources/cuda/kernels/spkron.cuh

‎cubool/include/cubool/cubool.h‎

Lines changed: 10 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -97,7 +97,9 @@ typedef enum cuBool_Hint {
9797
/** No duplicates in the build data */
9898
CUBOOL_HINT_NO_DUPLICATES = 256,
9999
/** Performs time measurement and logs elapsed operation time */
100-
CUBOOL_HINT_TIME_CHECK = 512
100+
CUBOOL_HINT_TIME_CHECK = 512,
101+
/** Transpose matrix before operation */
102+
CUBOOL_HINT_TRANSPOSE = 1024
101103
} cuBool_Hint;
102104

103105
/** Hit mask */
@@ -458,12 +460,18 @@ CUBOOL_EXPORT CUBOOL_API cuBool_Status cuBool_Matrix_Free(
458460
*
459461
* @return Error code on this operation
460462
*/
461-
CUBOOL_EXPORT CUBOOL_API cuBool_Status cuBool_Matrix_Reduce(
463+
CUBOOL_EXPORT CUBOOL_API cuBool_Status cuBool_Matrix_Reduce2(
462464
cuBool_Matrix result,
463465
cuBool_Matrix matrix,
464466
cuBool_Hints hints
465467
);
466468

469+
CUBOOL_EXPORT CUBOOL_API cuBool_Status cuBool_Matrix_Reduce(
470+
cuBool_Vector result,
471+
cuBool_Matrix matrix,
472+
cuBool_Hints hints
473+
);
474+
467475
/**
468476
* Performs result = left + right, where '+' is boolean semiring operation.
469477
*

‎cubool/sources/cuBool_Matrix_Reduce.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -25,16 +25,16 @@
2525
#include <cuBool_Common.hpp>
2626

2727
cuBool_Status cuBool_Matrix_Reduce(
28-
cuBool_Matrix result,
28+
cuBool_Vector result,
2929
cuBool_Matrix matrix,
3030
cuBool_Hints hints
3131
) {
3232
CUBOOL_BEGIN_BODY
3333
CUBOOL_VALIDATE_LIBRARY
3434
CUBOOL_ARG_NOT_NULL(result)
3535
CUBOOL_ARG_NOT_NULL(matrix)
36-
auto r = (cubool::Matrix*) result;
36+
auto r = (cubool::Vector*) result;
3737
auto m = (cubool::Matrix*) matrix;
38-
r->reduce(*m, hints & CUBOOL_HINT_TIME_CHECK);
38+
r->reduceMatrix(*m, hints & CUBOOL_HINT_TRANSPOSE,hints & CUBOOL_HINT_TIME_CHECK);
3939
CUBOOL_END_BODY
4040
}

‎cubool/sources/cuBool_Matrix_Reduce2.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_Reduce2(
28+
cuBool_Matrix result,
29+
cuBool_Matrix matrix,
30+
cuBool_Hints hints
31+
) {
32+
CUBOOL_BEGIN_BODY
33+
CUBOOL_VALIDATE_LIBRARY
34+
CUBOOL_ARG_NOT_NULL(result)
35+
CUBOOL_ARG_NOT_NULL(matrix)
36+
auto r = (cubool::Matrix*) result;
37+
auto m = (cubool::Matrix*) matrix;
38+
r->reduce(*m, hints & CUBOOL_HINT_TIME_CHECK);
39+
CUBOOL_END_BODY
40+
}

‎cubool/sources/cuda/cuda_matrix.cu‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -121,10 +121,6 @@ namespace cubool {
121121
return mMatrixImpl.m_vals;
122122
}
123123

124-
bool CudaMatrix::isStorageEmpty() const {
125-
return mMatrixImpl.is_zero_dim();
126-
}
127-
128124
bool CudaMatrix::isMatrixEmpty() const {
129125
return mMatrixImpl.m_vals == 0;
130126
}

‎cubool/sources/cuda/cuda_matrix.hpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -61,9 +61,10 @@ namespace cubool {
6161
index getNvals() const override;
6262

6363
private:
64+
friend class CudaVector;
65+
6466
void resizeStorageToDim() const;
6567
void clearAndResizeStorageToDim() const;
66-
bool isStorageEmpty() const;
6768
bool isMatrixEmpty() const;
6869
void transferToDevice(const std::vector<index> &rowOffsets, const std::vector<index> &colIndices) const;
6970
void transferFromDevice(std::vector<index> &rowOffsets, std::vector<index> &colIndices) const;

‎cubool/sources/cuda/cuda_vector.cu‎

Lines changed: 0 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -88,16 +88,6 @@ namespace cubool {
8888
result = getNvals();
8989
}
9090

91-
void CudaVector::reduceMatrix(const struct MatrixBase &matrix, bool transpose, bool checkTime) {
92-
RAISE_ERROR(NotImplemented, "This function is not implemented");
93-
94-
}
95-
96-
void CudaVector::eWiseAdd(const VectorBase &aBase, const VectorBase &bBase, bool checkTime) {
97-
RAISE_ERROR(NotImplemented, "This function is not implemented");
98-
99-
}
100-
10191
void CudaVector::multiplyVxM(const VectorBase &vBase, const struct MatrixBase &mBase, bool checkTime) {
10292
RAISE_ERROR(NotImplemented, "This function is not implemented");
10393

‎cubool/sources/cuda/cuda_vector_ewiseadd.cu‎

Lines changed: 47 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,47 @@
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 <cuda/cuda_vector.hpp>
26+
#include <cuda/kernels/spewiseadd.cuh>
27+
#include <core/error.hpp>
28+
#include <cassert>
29+
30+
namespace cubool {
31+
32+
void CudaVector::eWiseAdd(const VectorBase &aBase, const VectorBase &bBase, bool checkTime) {
33+
const auto* a = dynamic_cast<const CudaVector*>(&aBase);
34+
const auto* b = dynamic_cast<const CudaVector*>(&bBase);
35+
36+
CHECK_RAISE_ERROR(a != nullptr, InvalidArgument, "Provided vector does not belong to cuda vector class");
37+
CHECK_RAISE_ERROR(b != nullptr, InvalidArgument, "Provided vector does not belong to cuda vector class");
38+
39+
assert(a->getNrows() == b->getNrows());
40+
41+
kernels::SpVectorEWiseAdd<index, DeviceAlloc<index>> functor;
42+
auto result = functor(a->mVectorImpl, b->mVectorImpl);
43+
44+
mVectorImpl = std::move(result);
45+
}
46+
47+
}

‎cubool/sources/cuda/cuda_vector_reduce.cu‎

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,56 @@
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 <cuda/cuda_vector.hpp>
26+
#include <cuda/cuda_matrix.hpp>
27+
#include <cuda/kernels/spreduce.cuh>
28+
#include <core/error.hpp>
29+
#include <cassert>
30+
31+
namespace cubool {
32+
33+
void CudaVector::reduceMatrix(const struct MatrixBase &matrixBase, bool transpose, bool checkTime) {
34+
auto matrix = dynamic_cast<const CudaMatrix*>(&matrixBase);
35+
36+
CHECK_RAISE_ERROR(matrix != nullptr, InvalidArgument, "Provided matrix does not belongs to cuda matrix class");
37+
38+
if (transpose) {
39+
assert(matrix->getNcols() == this->getNrows());
40+
matrix->resizeStorageToDim();
41+
42+
// Reduce to row-vector
43+
kernels::SpVectorMatrixTransposedReduceFunctor<index, DeviceAlloc<index>> functor;
44+
mVectorImpl = std::move(functor(matrix->mMatrixImpl));
45+
}
46+
else {
47+
assert(matrix->getNrows() == this->getNrows());
48+
matrix->resizeStorageToDim();
49+
50+
// Reduce to column-vector
51+
kernels::SpVectorMatrixReduceFunctor<index, DeviceAlloc<index>> functor;
52+
mVectorImpl = std::move(functor(matrix->mMatrixImpl));
53+
}
54+
}
55+
56+
}

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

Lines changed: 70 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,70 @@
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+
#ifndef CUBOOL_SPEWISEADD_HPP
26+
#define CUBOOL_SPEWISEADD_HPP
27+
28+
#include <cuda/details/sp_vector.hpp>
29+
30+
namespace cubool {
31+
namespace kernels {
32+
33+
template <typename IndexType, typename AllocType>
34+
struct SpVectorEWiseAdd {
35+
template<typename T>
36+
using ContainerType = thrust::device_vector<T, typename AllocType::template rebind<T>::other>;
37+
using VectorType = details::SpVector<IndexType, AllocType>;
38+
39+
VectorType operator()(const VectorType& a, const VectorType& b) {
40+
auto aNvals = a.m_vals;
41+
auto bNvals = b.m_vals;
42+
auto worst = aNvals + bNvals;
43+
44+
// Allocate memory for the worst case scenario
45+
if (mCacheBuffer.size() < worst)
46+
mCacheBuffer.resize(aNvals + bNvals);
47+
48+
// Merge sorted arrays without duplicates
49+
auto out = thrust::set_union(a.m_rows_index.begin(), a.m_rows_index.end(),
50+
b.m_rows_index.begin(), b.m_rows_index.end(),
51+
mCacheBuffer.begin());
52+
53+
// Count result nvals count
54+
auto nvals = thrust::distance(mCacheBuffer.begin(), out);
55+
56+
// Fill the result buffer
57+
ContainerType<index> rowIndex(nvals);
58+
thrust::copy(mCacheBuffer.begin(), out, rowIndex.begin());
59+
60+
return VectorType(std::move(rowIndex), a.m_rows, nvals);
61+
}
62+
63+
private:
64+
ContainerType<index> mCacheBuffer;
65+
};
66+
67+
}
68+
}
69+
70+
#endif //CUBOOL_SPEWISEADD_HPP

0 commit comments

Comments
 (0)