Review update: rem code duplication, update YML file.

+ Remove code duplication in cuda kernels by moving common code to a .cuh file.
+ Update the artifacts uploading in the YML file to circumvent the GITLAB limits.

.gitlab-ci.yml

@@ -57,17 +57,18 @@ stages:
     paths:
       - "build/*/*/*/*/*/CMakeCache.txt"
       - "build/*/*/*/*/*/*.cmake"
-      - "build/*/*/*/*/*/core/test"
-      - "build/*/*/*/*/*/cuda/test"
-      - "build/*/*/*/*/*/omp/test"
-      - "build/*/*/*/*/*/reference/test"
+      - "build/*/*/*/*/*/core/test/*/[a-z_]*"
+      - "build/*/*/*/*/*/cuda/test/*/[a-z_]*"
+      - "build/*/*/*/*/*/omp/test/*/[a-z_]*"
+      - "build/*/*/*/*/*/reference/test/*/[a-z_]*"
       - "build/*/*/*/*/*/core/libginkgo*"
       - "build/*/*/*/*/*/cuda/libginkgo*"
       - "build/*/*/*/*/*/omp/libginkgo*"
       - "build/*/*/*/*/*/reference/libginkgo*"
       - "build/*/*/*/*/*/core/device_hooks/libginkgo*"
       - "build/*/*/*/*/*/*/CTestTestfile.cmake"
       - "build/*/*/*/*/*/*/*/CTestTestfile.cmake"
+      - "build/*/*/*/*/*/*/*/*/CTestTestfile.cmake"
   except:
       - schedules
 # build paths are of the form: build/<cuda_version>/<compiler>/<module(s)>/{debug,release}/{shared,static}/
@@ -140,7 +141,6 @@ build/cuda91/gcc/all/debug/static:
   image: localhost:5000/gko-cuda91-gnu6-llvm40
   variables:
     <<: *default_variables
-    BUILD_OMP: "ON"
     BUILD_CUDA: "ON"
     BUILD_TYPE: Debug
     EXTRA_CMAKE_FLAGS: *cuda_flags_static
@@ -187,7 +187,6 @@ build/cuda92/clang/all/debug/static:
     <<: *default_variables
     C_COMPILER: clang
     CXX_COMPILER: clang++
-    BUILD_OMP: "ON"
     BUILD_CUDA: "ON"
     BUILD_TYPE: Debug
     EXTRA_CMAKE_FLAGS: *cuda_flags_static

benchmark/spmv/cuda_linops.hpp

@@ -339,7 +339,7 @@ class CuspCsrEx
         // This function seems to require the pointer mode to be set to HOST.
         // Ginkgo use pointer mode DEVICE by default, so we change this
         // temporarily.
-        gko::cusparse_pointer_mode_guard pm_guard(handle);
+        gko::kernels::cuda::cusparse::pointer_mode_guard pm_guard(handle);
         gko::kernels::cuda::cusparse::spmv_buffersize<ValueType, IndexType>(
             handle, algmode_, trans_, this->get_size()[0], this->get_size()[1],
             csr_->get_num_stored_elements(), &alpha, this->get_descr(),

cuda/base/device_guard.hpp

@@ -30,6 +30,9 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 ******************************<GINKGO LICENSE>*******************************/
 
+#ifndef GKO_CUDA_BASE_DEVICE_GUARD_HPP_
+#define GKO_CUDA_BASE_DEVICE_GUARD_HPP_
+
 
 #include <cuda_runtime.h>
 
@@ -72,3 +75,5 @@ class device_guard {
 
 
 }  // namespace gko
+
+#endif

cuda/base/pointer_mode_guard.hpp

@@ -30,6 +30,9 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 ******************************<GINKGO LICENSE>*******************************/
 
+#ifndef GKO_CUDA_BASE_POINTER_MODE_GUARD_HPP_
+#define GKO_CUDA_BASE_POINTER_MODE_GUARD_HPP_
+
 
 #include <cublas_v2.h>
 #include <cuda_runtime.h>
@@ -40,28 +43,29 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 
 namespace gko {
+namespace kernels {
+namespace cuda {
+namespace cublas {
 
 
-class cublas_pointer_mode_guard {
+class pointer_mode_guard {
 public:
-    cublas_pointer_mode_guard(cublasHandle_t &handle)
+    pointer_mode_guard(cublasHandle_t &handle)
     {
         l_handle = &handle;
         GKO_ASSERT_NO_CUBLAS_ERRORS(
             cublasSetPointerMode(handle, CUBLAS_POINTER_MODE_HOST));
     }
 
-    cublas_pointer_mode_guard(cublas_pointer_mode_guard &other) = delete;
+    pointer_mode_guard(pointer_mode_guard &other) = delete;
 
-    cublas_pointer_mode_guard &operator=(
-        const cublas_pointer_mode_guard &other) = delete;
+    pointer_mode_guard &operator=(const pointer_mode_guard &other) = delete;
 
-    cublas_pointer_mode_guard(cublas_pointer_mode_guard &&other) = delete;
+    pointer_mode_guard(pointer_mode_guard &&other) = delete;
 
-    cublas_pointer_mode_guard const &operator=(
-        cublas_pointer_mode_guard &&other) = delete;
+    pointer_mode_guard const &operator=(pointer_mode_guard &&other) = delete;
 
-    ~cublas_pointer_mode_guard() noexcept(false)
+    ~pointer_mode_guard() noexcept(false)
     {
         /* Ignore the error during stack unwinding for this call */
         if (std::uncaught_exception()) {
@@ -77,26 +81,30 @@ class cublas_pointer_mode_guard {
 };
 
 
-class cusparse_pointer_mode_guard {
+}  // namespace cublas
+
+
+namespace cusparse {
+
+
+class pointer_mode_guard {
 public:
-    cusparse_pointer_mode_guard(cusparseHandle_t &handle)
+    pointer_mode_guard(cusparseHandle_t &handle)
     {
         l_handle = &handle;
         GKO_ASSERT_NO_CUSPARSE_ERRORS(
             cusparseSetPointerMode(handle, CUSPARSE_POINTER_MODE_HOST));
     }
 
-    cusparse_pointer_mode_guard(cusparse_pointer_mode_guard &other) = delete;
+    pointer_mode_guard(pointer_mode_guard &other) = delete;
 
-    cusparse_pointer_mode_guard &operator=(
-        const cusparse_pointer_mode_guard &other) = delete;
+    pointer_mode_guard &operator=(const pointer_mode_guard &other) = delete;
 
-    cusparse_pointer_mode_guard(cusparse_pointer_mode_guard &&other) = delete;
+    pointer_mode_guard(pointer_mode_guard &&other) = delete;
 
-    cusparse_pointer_mode_guard const &operator=(
-        cusparse_pointer_mode_guard &&other) = delete;
+    pointer_mode_guard const &operator=(pointer_mode_guard &&other) = delete;
 
-    ~cusparse_pointer_mode_guard() noexcept(false)
+    ~pointer_mode_guard() noexcept(false)
     {
         /* Ignore the error during stack unwinding for this call */
         if (std::uncaught_exception()) {
@@ -112,4 +120,9 @@ class cusparse_pointer_mode_guard {
 };
 
 
+}  // namespace cusparse
+}  // namespace cuda
+}  // namespace kernels
 }  // namespace gko
+
+#endif

cuda/matrix/csr_kernels.cu

@@ -697,14 +697,14 @@ void spmv(std::shared_ptr<const CudaExecutor> exec,
             auto handle = exec->get_cusparse_handle();
             auto descr = cusparse::create_mat_descr();
             {
-                cusparse_pointer_mode_guard pm_guard(handle);
+                cusparse::pointer_mode_guard pm_guard(handle);
                 auto row_ptrs = a->get_const_row_ptrs();
                 auto col_idxs = a->get_const_col_idxs();
                 auto alpha = one<ValueType>();
                 auto beta = zero<ValueType>();
-                if (b->get_stride() != 1 || c->get_stride() != 1)
+                if (b->get_stride() != 1 || c->get_stride() != 1) {
                     GKO_NOT_IMPLEMENTED;
-
+                }
                 cusparse::spmv(handle, CUSPARSE_OPERATION_NON_TRANSPOSE,
                                a->get_size()[0], a->get_size()[1],
                                a->get_num_stored_elements(), &alpha, descr,

cuda/matrix/dense_kernels.cu

@@ -72,7 +72,7 @@ void simple_apply(std::shared_ptr<const CudaExecutor> exec,
     if (cublas::is_supported<ValueType>::value) {
         auto handle = exec->get_cublas_handle();
         {
-            cublas_pointer_mode_guard pm_guard(handle);
+            cublas::pointer_mode_guard pm_guard(handle);
             auto alpha = one<ValueType>();
             auto beta = zero<ValueType>();
             cublas::gemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, c->get_size()[1],
@@ -937,7 +937,7 @@ void transpose(std::shared_ptr<const CudaExecutor> exec,
     if (cublas::is_supported<ValueType>::value) {
         auto handle = exec->get_cublas_handle();
         {
-            cublas_pointer_mode_guard pm_guard(handle);
+            cublas::pointer_mode_guard pm_guard(handle);
             auto alpha = one<ValueType>();
             auto beta = zero<ValueType>();
             cublas::geam(
@@ -963,7 +963,7 @@ void conj_transpose(std::shared_ptr<const CudaExecutor> exec,
     if (cublas::is_supported<ValueType>::value) {
         auto handle = exec->get_cublas_handle();
         {
-            cublas_pointer_mode_guard pm_guard(handle);
+            cublas::pointer_mode_guard pm_guard(handle);
             auto alpha = one<ValueType>();
             auto beta = zero<ValueType>();
             cublas::geam(