[Code] Add sub-vec extraction for cuda backend

Author: EgorOrachyov

cubool/sources/core/vector.cpp

@@ -85,7 +85,34 @@ namespace cubool {
     }
 
     void Vector::extractSubVector(const VectorBase &otherBase, index i, index nrows, bool checkTime) {
-        RAISE_ERROR(NotImplemented, "This function is not implemented");
+        const auto* other = dynamic_cast<const Vector*>(&otherBase);
+
+        CHECK_RAISE_ERROR(other != nullptr, InvalidArgument, "Passed vector does not belong to core vector class");
+
+        auto bI = i + nrows;
+
+        CHECK_RAISE_ERROR(nrows > 0, InvalidArgument, "Cannot extract sub-vector with zero dimension");
+        CHECK_RAISE_ERROR(bI <= other->getNrows(), InvalidArgument, "Provided sub-vector range must be within matrix bounds");
+        CHECK_RAISE_ERROR(nrows == this->getNrows(), InvalidArgument, "Result matrix has incompatible size for extracted sub-matrix range");
+
+        other->commitCache();
+        this->releaseCache(); // Values of this vector won't be used any more
+
+        if (checkTime) {
+            TIMER_ACTION(timer, mHnd->extractSubVector(*other->mHnd, i, nrows, false));
+
+            LogStream stream(*Library::getLogger());
+            stream << Logger::Level::Info
+                   << "Time: " << timer.getElapsedTimeMs() << " ms "
+                   << "Vector::extractSubVector: "
+                   << this->getDebugMarker() << " =subvector( "
+                   << i << ", shape=(" << nrows << ") "
+                   << other->getDebugMarker() << LogStream::cmt;
+
+            return;
+        }
+
+        mHnd->extractSubVector(*other->mHnd, i, nrows, false);
     }
 
     void Vector::clone(const VectorBase &otherBase) {

cubool/sources/cuda/cuda_vector.cu

@@ -25,6 +25,7 @@
 #include <cuda/cuda_vector.hpp>
 #include <core/error.hpp>
 #include <utils/data_utils.hpp>
+#include <limits>
 
 namespace cubool {
 
@@ -70,8 +71,61 @@ namespace cubool {
     }
 
     void CudaVector::extractSubVector(const VectorBase &otherBase, index i, index nrows, bool checkTime) {
-        RAISE_ERROR(NotImplemented, "This function is not implemented");
+        const auto* v = dynamic_cast<const CudaVector*>(&otherBase);
 
+        CHECK_RAISE_ERROR(v != nullptr, InvalidArgument, "Passed vector does not belong to cuda vector class");
+
+        assert(i + nrows <= v->getNrows());
+        assert(getNrows() == nrows);
+        assert(nrows > 0);
+
+        // If source is empty
+        if (v->getNvals() == 0) {
+            mVectorImpl = std::move(VectorImplType(nrows));
+            return;
+        }
+
+        // If source has too small values
+        index last = v->mVectorImpl.m_rows_index.back();
+        if (i > last) {
+            mVectorImpl = std::move(VectorImplType(nrows));
+            return;
+        }
+
+        auto& vec = v->mVectorImpl;
+
+        thrust::device_vector<index, DeviceAlloc<index>> region(2);
+        thrust::fill_n(region.begin(), 1, std::numeric_limits<index>::max());
+        thrust::fill_n(region.begin() + 1, 1, 0);
+
+        thrust::for_each(thrust::counting_iterator<index>(0), thrust::counting_iterator<index>(vec.m_vals),
+                         [first = region.data(), size = region.data() + 1,
+                          i = i, last = i + nrows, rowIndex = vec.m_rows_index.data()]
+                         __device__ (index id) {
+            auto rowId = rowIndex[id];
+
+            if (i <= rowId && rowId < last) {
+                atomicAdd(size.get(), 1);
+                atomicMin(first.get(), id);
+            }
+        });
+
+        index resultSize = region.back();
+
+        // If no values for result
+        if (resultSize == 0) {
+            mVectorImpl = std::move(VectorImplType(nrows));
+            return;
+        }
+
+        // Copy region to the result
+        index firstToCopy = region.front();
+
+        VectorImplType::container_type result(resultSize);
+        thrust::copy(vec.m_rows_index.begin() + firstToCopy, vec.m_rows_index.begin() + firstToCopy + resultSize, result.begin());
+
+        // Update this impl data
+        mVectorImpl = std::move(VectorImplType(std::move(result), nrows, resultSize));
     }
 
     void CudaVector::clone(const VectorBase &otherBase) {

cubool/sources/cuda/details/sp_vector.hpp

@@ -36,6 +36,8 @@ namespace cubool {
             typedef bool value_type;
             typedef IndexType index_type;
             typedef AllocType alloc_type;
+            typedef thrust::device_vector<IndexType, AllocType> container_type;
+
 
             SpVector()
                 : m_rows_index{}, m_rows{0}, m_vals{0} {}

cubool/tests/CMakeLists.txt

@@ -13,8 +13,8 @@ target_link_libraries(test_matrix_misc PUBLIC testing)
 add_executable(test_matrix_transpose test_matrix_transpose.cpp)
 target_link_libraries(test_matrix_transpose PUBLIC testing)
 
-add_executable(test_matrix_extract_sub_matrix test_matrix_extract_sub_matrix.cpp)
-target_link_libraries(test_matrix_extract_sub_matrix PUBLIC testing)
+add_executable(test_matrix_sub_matrix test_matrix_sub_matrix.cpp)
+target_link_libraries(test_matrix_sub_matrix PUBLIC testing)
 
 add_executable(test_matrix_reduce test_matrix_reduce.cpp)
 target_link_libraries(test_matrix_reduce PUBLIC testing)
@@ -42,3 +42,6 @@ target_link_libraries(test_vector_mxv PUBLIC testing)
 
 add_executable(test_vector_vxm test_vector_vxm.cpp)
 target_link_libraries(test_vector_vxm PUBLIC testing)
+
+add_executable(test_vector_sub_vector test_vector_sub_vector.cpp)
+target_link_libraries(test_vector_sub_vector PUBLIC testing)

cubool/tests/test_matrix_extract_sub_matrix.cpp renamed to cubool/tests/test_matrix_sub_matrix.cpp

@@ -0,0 +1,100 @@
+/**********************************************************************************/
+/* MIT License                                                                    */
+/*                                                                                */
+/* Copyright (c) 2020, 2021 JetBrains-Research                                    */
+/*                                                                                */
+/* Permission is hereby granted, free of charge, to any person obtaining a copy   */
+/* of this software and associated documentation files (the "Software"), to deal  */
+/* in the Software without restriction, including without limitation the rights   */
+/* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell      */
+/* copies of the Software, and to permit persons to whom the Software is          */
+/* furnished to do so, subject to the following conditions:                       */
+/*                                                                                */
+/* The above copyright notice and this permission notice shall be included in all */
+/* copies or substantial portions of the Software.                                */
+/*                                                                                */
+/* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR     */
+/* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,       */
+/* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE    */
+/* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER         */
+/* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,  */
+/* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE  */
+/* SOFTWARE.                                                                      */
+/**********************************************************************************/
+
+#include <testing/testing.hpp>
+
+void testVectorExtractSubVector(cuBool_Index m, cuBool_Index N, float density, cuBool_Hints flags) {
+    cuBool_Vector r, a;
+
+    auto ta = testing::Vector::generateSparse(m, density);
+
+    ASSERT_EQ(cuBool_Vector_New(&a, m), CUBOOL_STATUS_SUCCESS);
+    ASSERT_EQ(cuBool_Vector_Build(a, ta.index.data(), ta.nvals, CUBOOL_HINT_VALUES_SORTED), CUBOOL_STATUS_SUCCESS);
+
+    for (size_t i = 0; i < N; i++) {
+        auto k = m / N;
+
+        ASSERT_EQ(cuBool_Vector_New(&r, k), CUBOOL_STATUS_SUCCESS);
+        ASSERT_EQ(cuBool_Vector_ExtractSubVector(r, a, i * k, k, flags), CUBOOL_STATUS_SUCCESS);
+
+        auto tr = ta.subVector(i * k, k);
+
+        ASSERT_TRUE(tr.areEqual(r));
+        ASSERT_EQ(cuBool_Vector_Free(r), CUBOOL_STATUS_SUCCESS);
+    }
+
+    ASSERT_EQ(cuBool_Vector_Free(a), CUBOOL_STATUS_SUCCESS);
+}
+
+void testRun(cuBool_Index m, float step, cuBool_Hints setup) {
+    // Setup library
+    EXPECT_EQ(cuBool_Initialize(setup), CUBOOL_STATUS_SUCCESS);
+
+    auto N = 10;
+
+    for (size_t i = 0; i < N; i++) {
+        testVectorExtractSubVector(m, N, 0.1f + step * ((float) i), CUBOOL_HINT_NO);
+    }
+
+    // Finalize library
+    EXPECT_EQ(cuBool_Finalize(), CUBOOL_STATUS_SUCCESS);
+}
+
+TEST(cuBool_Matrix, SubVectorExtractSmall) {
+    cuBool_Index m = 10000;
+    float step = 0.05f;
+    testRun(m, step, CUBOOL_HINT_NO);
+}
+
+TEST(cuBool_Matrix, SubVectorExtractMedium) {
+    cuBool_Index m = 50000;
+    float step = 0.05f;
+    testRun(m, step, CUBOOL_HINT_NO);
+}
+
+TEST(cuBool_Matrix, SubVectorExtractLarge) {
+    cuBool_Index m = 100000;
+    float step = 0.05f;
+    testRun(m, step, CUBOOL_HINT_NO);
+}
+
+TEST(cuBool_Matrix, SubVectorExtractSmallFallback) {
+    cuBool_Index m = 10000;
+    float step = 0.05f;
+    testRun(m, step, CUBOOL_HINT_CPU_BACKEND);
+}
+
+TEST(cuBool_Matrix, SubVectorExtractMediumFallback) {
+    cuBool_Index m = 50000;
+    float step = 0.05f;
+    testRun(m, step, CUBOOL_HINT_CPU_BACKEND);
+}
+
+TEST(cuBool_Matrix, SubVectorExtractLargeFallback) {
+    cuBool_Index m = 100000;
+    float step = 0.05f;
+    testRun(m, step, CUBOOL_HINT_CPU_BACKEND);
+}
+
+CUBOOL_GTEST_MAIN

cubool/tests/test_vector_sub_vector.cpp

@@ -2,12 +2,12 @@
 # Invoke this script within build directory
 ./cubool/tests/test_library_api
 ./cubool/tests/test_matrix_ewiseadd
-./cubool/tests/test_matrix_extract_sub_matrix
 ./cubool/tests/test_matrix_kronecker
 ./cubool/tests/test_matrix_misc
 ./cubool/tests/test_matrix_mxm
 ./cubool/tests/test_matrix_reduce
 ./cubool/tests/test_matrix_setup
+./cubool/tests/test_matrix_sub_matrix
 ./cubool/tests/test_matrix_element
 ./cubool/tests/test_matrix_transpose
 ./cubool/tests/test_vector_misc

scripts/run_tests_all.sh

@@ -3,11 +3,11 @@
 ./cubool/tests/test_library_api
 ./cubool/tests/test_matrix_misc
 ./cubool/tests/test_matrix_ewiseadd --gtest_filter=*.*Fallback
-./cubool/tests/test_matrix_extract_sub_matrix --gtest_filter=*.*Fallback
 ./cubool/tests/test_matrix_kronecker --gtest_filter=*.*Fallback
 ./cubool/tests/test_matrix_mxm --gtest_filter=*.*Fallback
 ./cubool/tests/test_matrix_reduce --gtest_filter=*.*Fallback
 ./cubool/tests/test_matrix_setup --gtest_filter=*.*Fallback
+./cubool/tests/test_matrix_sub_matrix --gtest_filter=*.*Fallback
 ./cubool/tests/test_matrix_element --gtest_filter=*.*Fallback
 ./cubool/tests/test_matrix_transpose --gtest_filter=*.*Fallback
 ./cubool/tests/test_vector_misc