review update

Co-authored-by: Terry Cojean <terry.cojean@kit.edu>
Co-authored-by: Tobias Ribizel <ribizel@kit.edu>

.gitlab-ci.yml

@@ -336,26 +336,6 @@ build/cuda101/clang/all/release/static:
     - cuda
     - gpu
 
-build/cuda101/intel/cuda/debug/static:
-  <<: *default_build_with_test
-  image: localhost:5000/gko-cuda101-gnu8-llvm7-intel2019
-  variables:
-    <<: *default_variables
-    C_COMPILER: "icc"
-    CXX_COMPILER: "icpc"
-    BUILD_OMP: "ON"
-    BUILD_CUDA: "ON"
-    BUILD_TYPE: "Debug"
-    BUILD_SHARED_LIBS: "OFF"
-    CUDA_ARCH: 35
-  only:
-    variables:
-      - $RUN_CI_TAG
-  tags:
-    - private_ci
-    - cuda
-    - gpu
-
 # clang-cuda with cuda 10.1 and friends
 build/clang-cuda101/gcc/all/release/shared:
   <<: *default_build_with_test
@@ -458,14 +438,13 @@ build/cuda102/intel/cuda/debug/static:
     - gpu
 
 # cuda 11.0 and friends
-build/cuda110/gcc/all/debug/shared:
+build/cuda110/gcc/cuda/debug/shared:
   <<: *default_build_with_test
   image: localhost:5000/gko-cuda110-gnu9-llvm9-intel2020
   variables:
     <<: *default_variables
     BUILD_OMP: "ON"
     BUILD_CUDA: "ON"
-    BUILD_HIP: "OFF"
     BUILD_TYPE: "Debug"
     CUDA_ARCH: 35
   only:
@@ -476,7 +455,7 @@ build/cuda110/gcc/all/debug/shared:
     - cuda
     - gpu
 
-build/cuda110/clang/all/release/static:
+build/cuda110/clang/cuda/release/static:
   <<: *default_build_with_test
   image: localhost:5000/gko-cuda110-gnu9-llvm9-intel2020
   variables:
@@ -485,7 +464,6 @@ build/cuda110/clang/all/release/static:
     CXX_COMPILER: "clang++"
     BUILD_OMP: "ON"
     BUILD_CUDA: "ON"
-    BUILD_HIP: "OFF"
     BUILD_TYPE: "Release"
     BUILD_SHARED_LIBS: "OFF"
     CUDA_ARCH: 35

cuda/base/cusparse_bindings.hpp

@@ -214,10 +214,6 @@ GKO_BIND_CUSPARSE64_SPMV(ValueType, detail::not_implemented);
 #undef GKO_BIND_CUSPARSE64_SPMV
 
 
-#endif
-
-#if defined(CUDA_VERSION) && (CUDA_VERSION < 11000)
-
 #define GKO_BIND_CUSPARSE32_SPMM(ValueType, CusparseName)                     \
     inline void spmm(cusparseHandle_t handle, cusparseOperation_t transA,     \
                      int32 m, int32 n, int32 k, int32 nnz,                    \
@@ -265,6 +261,7 @@ GKO_BIND_CUSPARSE64_SPMM(ValueType, detail::not_implemented);
 #undef GKO_BIND_CUSPARSE32_SPMM
 #undef GKO_BIND_CUSPARSE64_SPMM
 
+
 #endif
 
 
@@ -351,6 +348,7 @@ GKO_BIND_CUSPARSE_SPMV_BUFFERSIZE(std::complex<double>);
 
 #undef GKO_BIND_CUSPARSE_SPMV_BUFFERSIZE
 
+
 #if defined(CUDA_VERSION) && (CUDA_VERSION < 11000)
 
 
@@ -379,7 +377,7 @@ GKO_BIND_CUSPARSE32_SPMV(ValueType, detail::not_implemented);
 #undef GKO_BIND_CUSPARSE32_SPMV
 
 
-#endif  // CUDA_VERSION < 11000
+#endif  // defined(CUDA_VERSION) && (CUDA_VERSION < 11000)
 
 
 #if defined(CUDA_VERSION) && (CUDA_VERSION < 11000)
@@ -627,6 +625,7 @@ GKO_BIND_CUSPARSE64_CSR2HYB(ValueType, detail::not_implemented);
 
 #endif
 
+
 #if defined(CUDA_VERSION) && (CUDA_VERSION < 11000)
 
 template <typename ValueType, typename IndexType>
@@ -734,6 +733,8 @@ GKO_BIND_CUSPARSE_TRANSPOSE32(float);
 GKO_BIND_CUSPARSE_TRANSPOSE32(double);
 GKO_BIND_CUSPARSE_TRANSPOSE32(std::complex<float>);
 GKO_BIND_CUSPARSE_TRANSPOSE32(std::complex<double>);
+
+
 #endif
 
 
@@ -768,7 +769,7 @@ inline void destroy(csrgemm2Info_t info)
 }
 
 
-#else
+#else  // CUDA_VERSION >= 11000
 
 
 inline cusparseSpGEMMDescr_t create_spgemm_descr()
@@ -1235,7 +1236,7 @@ GKO_BIND_CUSPARSE_GATHER(std::complex<double>, cusparseZgthr);
 #undef GKO_BIND_CUSPARSE_GATHER
 
 
-#else
+#else  // CUDA_VERSION >= 11000
 
 
 inline void gather(cusparseHandle_t handle, cusparseDnVecDescr_t in,

cuda/components/cooperative_groups.cuh

@@ -367,25 +367,17 @@ private:
 };
 
 
-#else
-
-
-// cuda11 put the default consturctor in protected function.
-template <typename Group>
-class enable_default_constructor : public Group {};
-
-
 #endif  // defined(CUDA_VERSION) && (CUDA_VERSION < 11000)
 
 
 }  // namespace detail
 
 
-// Implementing this as a using directive messes up with SFINAE for some reason,
-// probably a bug in NVCC. If it is a complete type, everything works fine.
 #if defined(CUDA_VERSION) && (CUDA_VERSION < 11000)
 
 
+// Implementing this as a using directive messes up with SFINAE for some reason,
+// probably a bug in NVCC. If it is a complete type, everything works fine.
 template <unsigned Size>
 struct thread_block_tile : detail::enable_extended_shuffle<
                                cooperative_groups::thread_block_tile<Size>> {
@@ -442,7 +434,7 @@ struct is_synchronizable_group_impl<cooperative_groups::thread_block_tile<Size>>
     : std::true_type {};
 
 
-#else
+#else  // CUDA_VERSION >= 11000
 
 
 // thread_block_tile is same as cuda11's

cuda/matrix/csr_kernels.cu

@@ -1111,7 +1111,7 @@ void conj_transpose(std::shared_ptr<const CudaExecutor> exec,
             gko::kernels::cuda::cuda_data_type<ValueType>();
         cusparseAction_t copyValues = CUSPARSE_ACTION_NUMERIC;
         cusparseIndexBase_t idxBase = CUSPARSE_INDEX_BASE_ZERO;
-        cusparseCsr2CscAlg_t alg = CUSPARSE_CSR2CSC_ALG2;
+        cusparseCsr2CscAlg_t alg = CUSPARSE_CSR2CSC_ALG1;
         size_type buffer_size = 0;
         cusparse::transpose_buffersize(
             exec->get_cusparse_handle(), orig->get_size()[0],

cuda/test/matrix/csr_kernels.cpp

@@ -416,6 +416,7 @@ TEST_F(Csr, TransposeIsEquivalentToRef)
 
     GKO_ASSERT_MTX_NEAR(static_cast<Mtx *>(d_trans.get()),
                         static_cast<Mtx *>(trans.get()), 0.0);
+    ASSERT_TRUE(static_cast<Mtx *>(d_trans.get())->is_sorted_by_column_index());
 }
 
 
@@ -428,6 +429,8 @@ TEST_F(Csr, ConjugateTransposeIsEquivalentToRef)
 
     GKO_ASSERT_MTX_NEAR(static_cast<ComplexMtx *>(d_trans.get()),
                         static_cast<ComplexMtx *>(trans.get()), 0.0);
+    ASSERT_TRUE(
+        static_cast<ComplexMtx *>(d_trans.get())->is_sorted_by_column_index());
 }
 
 

hip/test/matrix/csr_kernels.hip.cpp

@@ -401,6 +401,7 @@ TEST_F(Csr, TransposeIsEquivalentToRef)
 
     GKO_ASSERT_MTX_NEAR(static_cast<Mtx *>(d_trans.get()),
                         static_cast<Mtx *>(trans.get()), 0.0);
+    ASSERT_TRUE(static_cast<Mtx *>(d_trans.get())->is_sorted_by_column_index());
 }
 
 
@@ -482,18 +483,6 @@ TEST_F(Csr, MoveToSparsityCsrIsEquivalentToRef)
 }
 
 
-TEST_F(Csr, ConvertsEmptyToSellp)
-{
-    auto dempty_mtx = Mtx::create(hip);
-    auto dsellp_mtx = gko::matrix::Sellp<>::create(hip);
-
-    dempty_mtx->convert_to(dsellp_mtx.get());
-
-    ASSERT_EQ(hip->copy_val_to_host(dsellp_mtx->get_const_slice_sets()), 0);
-    ASSERT_FALSE(dsellp_mtx->get_size());
-}
-
-
 TEST_F(Csr, CalculateMaxNnzPerRowIsEquivalentToRef)
 {
     set_up_apply_data(std::make_shared<Mtx::sparselib>());
@@ -561,6 +550,18 @@ TEST_F(Csr, MoveToSellpIsEquivalentToRef)
 }
 
 
+TEST_F(Csr, ConvertsEmptyToSellp)
+{
+    auto dempty_mtx = Mtx::create(hip);
+    auto dsellp_mtx = gko::matrix::Sellp<>::create(hip);
+
+    dempty_mtx->convert_to(dsellp_mtx.get());
+
+    ASSERT_EQ(hip->copy_val_to_host(dsellp_mtx->get_const_slice_sets()), 0);
+    ASSERT_FALSE(dsellp_mtx->get_size());
+}
+
+
 TEST_F(Csr, CalculateTotalColsIsEquivalentToRef)
 {
     set_up_apply_data(std::make_shared<Mtx::sparselib>());
@@ -625,7 +626,7 @@ TEST_F(Csr, MoveToHybridIsEquivalentToRef)
 
 TEST_F(Csr, RecognizeSortedMatrixIsEquivalentToRef)
 {
-    set_up_apply_data(std::make_shared<Mtx::sparselib>());
+    set_up_apply_data(std::make_shared<Mtx::automatical>(hip));
     bool is_sorted_hip{};
     bool is_sorted_ref{};
 

omp/test/matrix/csr_kernels.cpp

@@ -299,6 +299,7 @@ TEST_F(Csr, TransposeIsEquivalentToRef)
 
     GKO_ASSERT_MTX_NEAR(static_cast<Mtx *>(d_trans.get()),
                         static_cast<Mtx *>(trans.get()), 0.0);
+    ASSERT_TRUE(static_cast<Mtx *>(d_trans.get())->is_sorted_by_column_index());
 }
 
 
@@ -311,6 +312,8 @@ TEST_F(Csr, ConjugateTransposeIsEquivalentToRef)
 
     GKO_ASSERT_MTX_NEAR(static_cast<ComplexMtx *>(d_trans.get()),
                         static_cast<ComplexMtx *>(trans.get()), 0.0);
+    ASSERT_TRUE(
+        static_cast<ComplexMtx *>(d_trans.get())->is_sorted_by_column_index());
 }