[Code] Add SpGEMV first impl for cuda backend

Author: EgorOrachyov

cubool/CMakeLists.txt

@@ -128,6 +128,7 @@ if (CUBOOL_WITH_CUDA)
         sources/cuda/cuda_matrix_extract_sub_matrix.cu
         sources/cuda/cuda_vector.hpp
         sources/cuda/cuda_vector.cu
+        sources/cuda/cuda_vector_mxv.cu
         sources/cuda/cuda_vector_ewiseadd.cu
         sources/cuda/cuda_vector_reduce.cu
         sources/cuda/details/meta.hpp
@@ -136,6 +137,7 @@ if (CUBOOL_WITH_CUDA)
         sources/cuda/details/device_allocator.cuh
         sources/cuda/kernels/slow_sort.cuh
         sources/cuda/kernels/bin_search.cuh
+        sources/cuda/kernels/spgemv.cuh
         sources/cuda/kernels/spewiseadd.cuh
         sources/cuda/kernels/sptranspose.cuh
         sources/cuda/kernels/sptranspose2.cuh

cubool/sources/core/library.cpp

@@ -240,6 +240,10 @@ namespace cubool {
 
     void Library::handleError(const std::exception& error) {
         mLogger->log(Logger::Level::Error, error.what());
+
+#ifdef CUBOOL_DEBUG
+        std::cerr << error.what() << std::endl;
+#endif
     }
 
     void Library::queryCapabilities(cuBool_DeviceCaps &caps) {

cubool/sources/cuda/cuda_vector.cu

@@ -93,11 +93,6 @@ namespace cubool {
 
     }
 
-    void CudaVector::multiplyMxV(const struct MatrixBase &mBase, const VectorBase &vBase, bool checkTime) {
-        RAISE_ERROR(NotImplemented, "This function is not implemented");
-
-    }
-
     index CudaVector::getNrows() const {
         return mVectorImpl.m_rows;
     }

cubool/sources/cuda/cuda_vector_mxv.cu

@@ -0,0 +1,51 @@
+/**********************************************************************************/
+/* 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 <cuda/cuda_vector.hpp>
+#include <cuda/cuda_matrix.hpp>
+#include <cuda/kernels/spgemv.cuh>
+#include <core/error.hpp>
+#include <cassert>
+
+namespace cubool {
+
+    void CudaVector::multiplyMxV(const struct MatrixBase &mBase, const VectorBase &vBase, bool checkTime) {
+        const auto* m = dynamic_cast<const CudaMatrix*>(&mBase);
+        const auto* v = dynamic_cast<const CudaVector*>(&vBase);
+
+        CHECK_RAISE_ERROR(m != nullptr, InvalidArgument, "Provided matrix does not belong to cuda matrix class");
+        CHECK_RAISE_ERROR(v != nullptr, InvalidArgument, "Provided vector does not belong to cuda vector class");
+
+        assert(m->getNcols() == v->getNrows());
+        assert(m->getNrows() == this->getNrows());
+
+        m->resizeStorageToDim();
+
+        kernels::SpGEMV<index, DeviceAlloc<index>> functor;
+        auto result = functor(m->mMatrixImpl, v->mVectorImpl);
+
+        mVectorImpl = std::move(result);
+    }
+
+}

cubool/sources/cuda/details/meta.hpp

@@ -25,59 +25,64 @@
 #ifndef CUBOOL_META_HPP
 #define CUBOOL_META_HPP
 
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
 #include <cstddef>
 
 namespace cubool {
 
-    template <size_t workersCount, size_t blockSize>
-    struct BinConfig {
-        static const size_t mWorkersCount = workersCount;
-        static const size_t mBlockSize = blockSize;
-    };
+    template <typename Config>
+    struct StreamsWrapper {
+        StreamsWrapper() {
+            for (auto& s: streams)
+                cudaStreamCreate(&s);
+        }
 
-    template <size_t minBorder, size_t maxBorder>
-    struct BorderConfig {
-        static const size_t mMinBorder = minBorder;
-        static const size_t mMaxBorder = maxBorder;
-    };
+        ~StreamsWrapper() {
+            for (auto& s: streams)
+                cudaStreamDestroy(s);
+        }
 
-    template <typename ... BinsConfig>
-    class ComputeConfig {
-    public:
+        cudaStream_t streams[Config::binsCount()] = {};
+    };
 
-        template<typename T, typename C>
-        void exec(C& context) {
-            ExecImpl<void, T, BinsConfig...>::exec(0, context);
-        }
+    template<size_t BlocksSize, size_t Max, size_t Min, size_t Id>
+    struct Bin {
+        static constexpr size_t blockSize = BlocksSize;
+        static constexpr size_t min = Max;
+        static constexpr size_t max = Min;
+        static constexpr size_t id = Id;
+    };
 
-    private:
 
-        template<typename D, typename T, typename C, typename ... Bins>
-        struct ExecImpl {
-            static void exec(size_t index, C& context) { }
-        };
+    template <typename ... Bins>
+    struct Config {
+    public:
 
-        template<typename D, typename T, typename C, typename Bin>
-        struct ExecImpl<D, T, C, Bin> {
-            static void exec(size_t index, C& context) {
-                using Timpl = typename T::template function<Bin>;
+        static __host__ __device__ size_t selectBin(size_t rowSize) {
+            static constexpr size_t mins[] = { Bins::min... };
+            static constexpr size_t maxs[] = { Bins::max... };
 
-                Timpl timpl;
-                timpl(index, context);
+            for (size_t i = 0; i < binsCount(); i++) {
+                if (mins[i] <= rowSize && rowSize <= maxs[i])
+                    return i;
             }
-        };
 
-        template<typename D, typename T, typename C, typename Bin, typename ... Bins>
-        struct ExecImpl<D, T, C, Bin, Bins...> {
-            static void exec(size_t index, C& context) {
-                using Timpl = typename T::template function<Bin>;
+            return unusedBinId();
+        }
 
-                Timpl timpl;
-                timpl(index, context);
+        static __host__ __device__ constexpr size_t binBlockSize(size_t id) {
+            constexpr size_t blockSizes[] = { Bins::blockSize... };
+            return blockSizes[id];
+        }
 
-                ExecImpl<D, T, Bins...>::exec(index + 1, context);
-            }
-        };
+        static __host__ __device__ constexpr size_t binsCount() {
+            return sizeof...(Bins);
+        }
+
+        static __host__ __device__ constexpr size_t unusedBinId() {
+            return binsCount() + 1;
+        }
 
     };