add cuda 13.3 CI job and workaround for typename decltype issue in cuda 13.3

.gitlab-ci.yml

@@ -219,6 +219,19 @@ build/cuda131/nompi/gcc/cuda/release/shared:
     BUILD_TYPE: "Release"
     MODULE_LOAD: "cmake/3.30.8 cuda/13.1.1 gcc/14.3.0"
 
+build/cuda133/nompi/gcc/cuda/release/shared:
+  extends:
+    - .build_and_test_tum_template
+    - .default_variables
+    - .full_test_condition
+    - .use_tum-nvidia
+  variables:
+    BUILD_CUDA: "ON"
+    ENABLE_HALF: "ON"
+    ENABLE_BFLOAT16: "ON"
+    BUILD_TYPE: "Release"
+    MODULE_LOAD: "cmake/3.31.6 cuda/13.3.0 gcc/14.3.0"
+
 build/amd/openmpi/gcc/rocm620/debug/shared:
   extends:
     - .build_and_test_tum_template
@@ -261,7 +274,7 @@ build/amd/nompi/gcc/rocm644/debug/static:
     BUILD_SHARED_LIBS: "OFF"
     MODULE_LOAD: "cmake/3.29.6 rocm/6.4.4 gcc/14.3.0"
 
-build/amd/nompi/gcc/rocm710/release/shared:
+build/amd/nompi/gcc/rocm711/release/shared:
   extends:
     - .build_and_test_tum_template
     - .default_variables
@@ -272,7 +285,7 @@ build/amd/nompi/gcc/rocm710/release/shared:
     BUILD_MPI: "OFF"
     BUILD_OMP: "OFF"
     BUILD_TYPE: "Release"
-    MODULE_LOAD: "cmake/3.29.6 rocm/7.1.0 gcc/14.3.0"
+    MODULE_LOAD: "cmake/3.29.6 rocm/7.1.1 gcc/14.3.0"
 
 build/amd/nompi/gcc/rocm720/release/shared:
   extends:

common/cuda_hip/components/prefix_sum_kernels.cpp

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -59,7 +59,6 @@ template <typename IndexType>
 void prefix_sum_nonnegative(std::shared_ptr<const DefaultExecutor> exec,
                             IndexType* counts, size_type num_entries)
 {
-    constexpr auto max = std::numeric_limits<IndexType>::max();
     thrust::exclusive_scan(thrust_policy(exec), counts, counts + num_entries,
                            counts, IndexType{}, overflowing_sum<IndexType>{});
     if (num_entries > 0 &&

common/cuda_hip/factorization/lu_kernels.cpp

@@ -107,7 +107,6 @@ __global__ __launch_bounds__(default_block_size) void factorize(
     const auto lane = warp.thread_rank();
     const auto row_begin = row_ptrs[row];
     const auto row_diag = diag_idxs[row];
-    const auto row_end = row_ptrs[row + 1];
     gko::matrix::csr::device_sparsity_lookup<IndexType> lookup{
         row_ptrs, cols,      storage_offsets,
         storage,  row_descs, static_cast<size_type>(row)};

common/cuda_hip/matrix/csr_kernels.template.cpp

@@ -1741,7 +1741,6 @@ void compute_submatrix(
     auto col_offset = col_span.begin;
     auto num_rows = result.size[0];
     auto num_cols = result.size[1];
-    auto row_ptrs = source.row_ptrs;
     auto grid_dim = ceildiv(num_rows, default_block_size);
     if (grid_dim > 0) {
         kernel::compute_submatrix_idxs_and_vals<<<grid_dim, default_block_size,
@@ -2683,7 +2682,6 @@ void advanced_spgemm(std::shared_ptr<const DefaultExecutor> exec,
     auto b_row_ptrs = b.row_ptrs;
     auto b_col_idxs = b.col_idxs;
     auto vbeta = exec->copy_val_to_host(beta.values);
-    auto d_nnz = IndexType(d.num_stored_elements);
     auto d_vals = d.values;
     auto d_row_ptrs = d.row_ptrs;
     auto d_col_idxs = d.col_idxs;

common/cuda_hip/matrix/dense_kernels.cpp

@@ -612,7 +612,6 @@ void convert_to_sellp(std::shared_ptr<const DefaultExecutor> exec,
     auto slice_sets = result.slice_sets;
 
     const auto slice_size = result.slice_size;
-    const auto stride_factor = result.stride_factor;
 
     auto grid_dim = ceildiv(num_rows, default_block_size / config::warp_size);
     if (grid_dim > 0) {

common/cuda_hip/matrix/ell_kernels.cpp

@@ -180,7 +180,6 @@ __global__ __launch_bounds__(default_block_size) void spmv(
 {
     using arithmetic_type = typename a_accessor::arithmetic_type;
     const auto alpha_val = alpha(0);
-    const OutputValueType beta_val = beta[0];
     if constexpr (atomic) {
         // Because the atomic operation changes the values of c during
         // computation, it can not directly do alpha * a * b + beta * c
@@ -194,6 +193,7 @@ __global__ __launch_bounds__(default_block_size) void spmv(
                 return static_cast<OutputValueType>(alpha_val * x);
             });
     } else {
+        const OutputValueType beta_val = beta[0];
         if (is_zero(beta_val)) {
             spmv_kernel<num_thread_per_worker, atomic>(
                 num_rows, num_worker_per_row, val, col, stride,

common/unified/matrix/sellp_kernels.cpp

@@ -168,8 +168,6 @@ void convert_to_csr(
             const auto slice = row / slice_size;
             const auto local_row = row % slice_size;
             const auto slice_begin = slice_sets[slice];
-            const auto slice_end = slice_sets[slice + 1];
-            const auto slice_length = slice_end - slice_begin;
             auto in_idx = slice_begin * slice_size + local_row;
             for (auto i = row_begin; i < row_end; i++) {
                 out_cols[i] = cols[in_idx];

core/base/dense_cache.cpp

@@ -106,7 +106,8 @@ ScalarCacheAccessor::get_scalars(const ScalarCache& cache)
 
 ScalarCache::ScalarCache(std::shared_ptr<const Executor> executor,
                          double scalar_value)
-    : exec(std::move(executor)), value(scalar_value){};
+    : exec(std::move(executor)), value(scalar_value)
+{}
 
 ScalarCache::ScalarCache(const ScalarCache& other) { *this = other; }
 

core/matrix/dense.cpp

@@ -2032,15 +2032,15 @@ auto Dense<ValueType>::get_device_view() -> device_view
 {
     return device_view{this->get_size(), this->get_stride(),
                        this->get_values()};
-};
+}
 
 
 template <typename ValueType>
 auto Dense<ValueType>::get_const_device_view() const -> const_device_view
 {
     return const_device_view{this->get_size(), this->get_stride(),
                              this->get_const_values()};
-};
+}
 
 
 template <typename ValueType>

include/ginkgo/core/base/abstract_factory.hpp

@@ -18,6 +18,39 @@
  * @ingroup gko
  */
 namespace gko {
+namespace detail {
+
+/**
+ * workaround for nvcc 13.3 complains `typename decltype(_name)::element_type`
+ * needs `typename` before `decltype`.
+ *
+ * TODO: remove workaround for nvcc 13.3. see
+ * https://github.com/ginkgo-project/ginkgo/pull/2076
+ *
+ * @note can not use `using element_type_t = typename T::element_type`.
+ * cuda 13.1 has difficulty to resolve it in Isai.
+ */
+template <typename T>
+struct element_type {
+    using type = typename T::element_type;
+};
+
+template <typename T>
+using element_type_t = typename element_type<T>::type;
+
+
+/**
+ * workaround for nvcc 13.3 complains `typename decltype(_name)::element_type`
+ * needs `typename` before `decltype`.
+ *
+ * TODO: remove workaround for nvcc 13.3. see
+ * https://github.com/ginkgo-project/ginkgo/pull/2076
+ */
+template <typename T>
+using value_element_type_t = typename T::value_type::element_type;
+
+
+}  // namespace detail
 
 
 /**
@@ -514,7 +547,7 @@ class deferred_factory_parameter {
     _name{};                                                                   \
                                                                                \
 private:                                                                       \
-    using _name##_type = typename decltype(_name)::element_type;               \
+    using _name##_type = ::gko::detail::element_type_t<decltype(_name)>;       \
                                                                                \
 public:                                                                        \
     auto with_##_name(::gko::deferred_factory_parameter<_name##_type> factory) \
@@ -552,7 +585,7 @@ public:                                                                        \
     _name{};                                                                   \
                                                                                \
 private:                                                                       \
-    using _name##_type = typename decltype(_name)::value_type::element_type;   \
+    using _name##_type = ::gko::detail::value_element_type_t<decltype(_name)>; \
                                                                                \
 public:                                                                        \
     template <typename... Args,                                                \

omp/preconditioner/sor_kernels.cpp

@@ -28,7 +28,7 @@ void initialize_weighted_l(
         factorization::helpers::triangular_mtx_closure(
             [inv_weight](auto val) { return val * inv_weight; },
             [](auto val) { return val; }));
-};
+}
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
     GKO_DECLARE_SOR_INITIALIZE_WEIGHTED_L);

test/components/bitvector.cpp

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -35,7 +35,6 @@ class Bitvector : public CommonTestFixture {
     using bitvector = gko::bitvector<index_type>;
     using device_bitvector = gko::device_bitvector<index_type>;
     using storage_type = typename bitvector::storage_type;
-    constexpr static auto block_size = bitvector::block_size;
 
     Bitvector()
         : rng{67193}, sizes{0,    1,    2,    16,    31,    32,  33,
@@ -173,7 +172,6 @@ TYPED_TEST(Bitvector, AccessIsEquivalentToRef)
 {
     using index_type = typename TestFixture::index_type;
     using storage_type = typename TestFixture::storage_type;
-    constexpr auto block_size = TestFixture::block_size;
     for (auto size : this->sizes) {
         SCOPED_TRACE(size);
         for (auto num_values :