diff --git a/.github/workflows/github-to-codehub-push-sync.yml b/.github/workflows/github-to-codehub-push-sync.yml deleted file mode 100644 index f486b34aae7..00000000000 --- a/.github/workflows/github-to-codehub-push-sync.yml +++ /dev/null @@ -1,24 +0,0 @@ -name: Codehub-sync - -on: - push: - branches: - - 'main' - - 'develop' - - 'release/**' - tags: - - 'v[0-9]+.[0-9]+.[0-9]+' - -jobs: - sync: - runs-on: ubuntu-latest - name: CASTIEL2 Codehub Sync - steps: - - uses: actions/checkout@v4 - with: - fetch-depth: 0 - - uses: cniethammer/git-repo-sync@cb1067d118c0c209bcded585b5c78e581bb0dfec - with: - target-url: 'https://codehub.hlrs.de/coes/microcard2/ginkgo.git' - target-username: ${{ secrets.CODEHUB_TOKEN_NAME }} - target-token: ${{ secrets.CODEHUB_TOKEN }} diff --git a/.github/workflows/intel.yml b/.github/workflows/intel.yml deleted file mode 100644 index c493b380c10..00000000000 --- a/.github/workflows/intel.yml +++ /dev/null @@ -1,57 +0,0 @@ -name: Intel-build - -on: - push: - branches: - - 'main' - - 'master' - - 'develop' - - 'release/**' - tags: - - '**' - pull_request: - types: [opened,synchronize] - paths-ignore: - - 'doc/**' - -concurrency: - group: ${{ github.workflow }}-${{ (github.head_ref && github.ref) || github.run_id }} - cancel-in-progress: true - -jobs: - intel: - strategy: - fail-fast: false - matrix: - config: - - {compiler: "dpcpp", build_type: "Release", name: "intel/dpcpp/release/shared", mixed: "ON", half: "ON", bfloat16: "OFF"} - - {compiler: "icpx", build_type: "Release", name: "intel/icpx/release/shared", mixed: "OFF", half: "ON", bfloat16: "OFF"} - - {compiler: "icpx", build_type: "Release", name: "intel/icpx/release/shared-bfloat16", mixed: "ON", half: "OFF", bfloat16: "ON"} - - {compiler: "icpx", build_type: "Release", name: "intel/icpx/release/shared-16bit", mixed: "ON", half: "ON", bfloat16: "ON"} - name: ${{ matrix.config.name }} - runs-on: [gpu_intel] - - steps: - - name: Checkout the latest code (shallow clone) - uses: actions/checkout@v4 - - - name: configure - run: | - source /etc/profile - module load intel-oneapi-compilers intel-oneapi-dpl intel-oneapi-mkl cmake - mkdir build - cd build - cmake .. -DCMAKE_INSTALL_PREFIX=install_ginkgo -DCMAKE_CXX_FLAGS="-Wpedantic -ffp-model=precise" -DCMAKE_CXX_COMPILER=${{ matrix.config.compiler }} -DCMAKE_BUILD_TYPE=${{ matrix.config.build_type }} -DGINKGO_MIXED_PRECISION=${{ matrix.config.mixed }} -DGINKGO_BUILD_CUDA=OFF -DGINKGO_BUILD_HIP=OFF -DGINKGO_BUILD_MPI=OFF -DGINKGO_DPCPP_SINGLE_MODE=ON -DGINKGO_ENABLE_HALF=${{ matrix.config.half }} -DGINKGO_ENABLE_BFLOAT16=${{ matrix.config.bfloat16 }} - make -j8 - ONEAPI_DEVICE_SELECTOR=level_zero:gpu ctest -j10 --output-on-failure - - - name: install - run: | - source /etc/profile - module load intel-oneapi-compilers intel-oneapi-dpl intel-oneapi-mkl cmake - cd build - SYCL_DEVICE_FILTER=level_zero:gpu make install - export GINKGO_PATH="$(pwd)/install_ginkgo/lib" - export LIBRARY_PATH=${ICL_INTEL_TBB_ROOT}/lib64:${GINKGO_PATH}:$LIBRARY_PATH - export LD_LIBRARY_PATH=${ICL_INTEL_TBB_ROOT}/lib64:${GINKGO_PATH}:$LD_LIBRARY_PATH - SYCL_DEVICE_FILTER=level_zero:gpu make test_install diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml index ab672b76045..1f9680191ad 100644 --- a/.gitlab-ci.yml +++ b/.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: diff --git a/common/cuda_hip/components/prefix_sum_kernels.cpp b/common/cuda_hip/components/prefix_sum_kernels.cpp index ebf102a7181..d5e3718f49c 100644 --- a/common/cuda_hip/components/prefix_sum_kernels.cpp +++ b/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 void prefix_sum_nonnegative(std::shared_ptr exec, IndexType* counts, size_type num_entries) { - constexpr auto max = std::numeric_limits::max(); thrust::exclusive_scan(thrust_policy(exec), counts, counts + num_entries, counts, IndexType{}, overflowing_sum{}); if (num_entries > 0 && diff --git a/common/cuda_hip/factorization/lu_kernels.cpp b/common/cuda_hip/factorization/lu_kernels.cpp index 2e955582af2..6b9cbd65228 100644 --- a/common/cuda_hip/factorization/lu_kernels.cpp +++ b/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 lookup{ row_ptrs, cols, storage_offsets, storage, row_descs, static_cast(row)}; diff --git a/common/cuda_hip/matrix/csr_kernels.template.cpp b/common/cuda_hip/matrix/csr_kernels.template.cpp index 8d3f96f4388..c77783c96ce 100644 --- a/common/cuda_hip/matrix/csr_kernels.template.cpp +++ b/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<< 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; diff --git a/common/cuda_hip/matrix/dense_kernels.cpp b/common/cuda_hip/matrix/dense_kernels.cpp index 19f7eb732b6..6ad6fb6975f 100644 --- a/common/cuda_hip/matrix/dense_kernels.cpp +++ b/common/cuda_hip/matrix/dense_kernels.cpp @@ -612,7 +612,6 @@ void convert_to_sellp(std::shared_ptr 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) { diff --git a/common/cuda_hip/matrix/ell_kernels.cpp b/common/cuda_hip/matrix/ell_kernels.cpp index 5634bc80c6e..65a45ccbd2c 100644 --- a/common/cuda_hip/matrix/ell_kernels.cpp +++ b/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(alpha_val * x); }); } else { + const OutputValueType beta_val = beta[0]; if (is_zero(beta_val)) { spmv_kernel( num_rows, num_worker_per_row, val, col, stride, diff --git a/common/unified/matrix/sellp_kernels.cpp b/common/unified/matrix/sellp_kernels.cpp index 7132ea13d12..f7884564936 100644 --- a/common/unified/matrix/sellp_kernels.cpp +++ b/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]; diff --git a/core/base/dense_cache.cpp b/core/base/dense_cache.cpp index 770a9037333..61bee55e9c9 100644 --- a/core/base/dense_cache.cpp +++ b/core/base/dense_cache.cpp @@ -106,7 +106,8 @@ ScalarCacheAccessor::get_scalars(const ScalarCache& cache) ScalarCache::ScalarCache(std::shared_ptr 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; } diff --git a/core/matrix/dense.cpp b/core/matrix/dense.cpp index 2e1f310e73e..4116ffa41ca 100644 --- a/core/matrix/dense.cpp +++ b/core/matrix/dense.cpp @@ -2032,7 +2032,7 @@ auto Dense::get_device_view() -> device_view { return device_view{this->get_size(), this->get_stride(), this->get_values()}; -}; +} template @@ -2040,7 +2040,7 @@ auto Dense::get_const_device_view() const -> const_device_view { return const_device_view{this->get_size(), this->get_stride(), this->get_const_values()}; -}; +} template diff --git a/include/ginkgo/core/base/abstract_factory.hpp b/include/ginkgo/core/base/abstract_factory.hpp index c6c863bf316..3a4b4d2293b 100644 --- a/include/ginkgo/core/base/abstract_factory.hpp +++ b/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 +struct element_type { + using type = typename T::element_type; +}; + +template +using element_type_t = typename element_type::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 +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; \ \ 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; \ \ public: \ template ; using device_bitvector = gko::device_bitvector; 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 :