Run Intel CI jobs on cluster

.gitlab-ci.yml

@@ -339,12 +339,11 @@ build/nocuda-nomixed/nompi/clang/omp/debug/static:
     MIXED_PRECISION: "OFF"
 
 # spack oneapi 2023.1
-build/icpx20231/igpu/release/shared:
+build/icpx20231/gpu/release/shared:
   extends:
     - .build_and_test_template
     - .default_variables
     - .quick_test_condition
-    - .disable_job_condition
     - .use_gko-oneapi20231-igpu
   variables:
     CXX_COMPILER: "icpx"
@@ -373,19 +372,19 @@ build/icpx20231/igpu/release/shared:
 #     ONEAPI_DEVICE_SELECTOR: "level_zero:gpu"
 
 # It gives two available backends of GPU on tests
-build/icpx/igpu/release/static:
+build/dpcpp/gpu/release/shared:
   extends:
     - .build_and_test_template
     - .default_variables
     - .full_test_condition
-    - .disable_job_condition
     - .use_gko-oneapi-igpu
   variables:
     CXX_COMPILER: "dpcpp"
     CXX_FLAGS: "-Wpedantic -ffp-model=precise"
     BUILD_SYCL: "ON"
     BUILD_TYPE: "Release"
-    BUILD_SHARED_LIBS: "OF"
+    # static builds take too long
+    BUILD_SHARED_LIBS: "ON"
     DPCPP_SINGLE_MODE: "ON"
     ONEAPI_DEVICE_SELECTOR: "*:gpu"
     BUILD_HWLOC: "OFF"

.gitlab/image.yml

@@ -50,17 +50,14 @@
 .use_gko-oneapi-igpu:
   image: ginkgohub/oneapi:latest
   tags:
-    - private_ci
-    - intel-igpu
+    - intel-gpus
 
 .use_gko-oneapi20231-igpu:
   image: ginkgohub/spack-oneapi:20231-openmpi
   tags:
-    - private_ci
-    - intel-igpu
+    - intel-gpus
 
 .use_gko-oneapi-dgpu:
   image: ginkgohub/oneapi:latest
   tags:
-    - private_ci
-    - intel-dgpu
+    - intel-gpus

dpcpp/base/executor.dp.cpp

@@ -279,7 +279,11 @@ void DpcppExecutor::set_device_property(dpcpp_queue_property property)
     // `wait()` would be needed after every call to a DPC++ function or kernel.
     // For example, without `in_order`, doing a copy, a kernel, and a copy, will
     // not necessarily happen in that order by default, which we need to avoid.
-    auto* queue = new sycl::queue{device, detail::get_property_list(property)};
+    // We need to create the context for each device. Otherwise, we get -999
+    // Unknown PI error after second device.
+    // Ref: https://github.com/intel/llvm/issues/10982
+    auto* queue = new sycl::queue{sycl::context(device), device,
+                                  detail::get_property_list(property)};
     queue_ = std::move(queue_manager<sycl::queue>{queue, detail::delete_queue});
 }
 

dpcpp/preconditioner/sor_kernels.dp.cpp

@@ -31,14 +31,18 @@ void initialize_weighted_l(
                         1, 1};
 
     auto inv_weight = one(weight) / weight;
+    const auto in_row_ptrs = system_matrix->get_const_row_ptrs();
+    const auto in_col_idxs = system_matrix->get_const_col_idxs();
+    const auto in_values = system_matrix->get_const_values();
+    const auto l_row_ptrs = l_mtx->get_const_row_ptrs();
+    const auto l_col_idxs = l_mtx->get_col_idxs();
+    const auto l_values = l_mtx->get_values();
 
     exec->get_queue()->parallel_for(
         sycl_nd_range(grid_dim, block_size), [=](sycl::nd_item<3> item_ct1) {
             factorization::helpers::initialize_l(
-                num_rows, system_matrix->get_const_row_ptrs(),
-                system_matrix->get_const_col_idxs(),
-                system_matrix->get_const_values(), l_mtx->get_const_row_ptrs(),
-                l_mtx->get_col_idxs(), l_mtx->get_values(),
+                num_rows, in_row_ptrs, in_col_idxs, in_values, l_row_ptrs,
+                l_col_idxs, l_values,
                 factorization::helpers::triangular_mtx_closure(
                     [inv_weight](auto val) { return val * inv_weight; },
                     factorization::helpers::identity{}),
@@ -67,15 +71,21 @@ void initialize_weighted_l_u(
     auto inv_two_minus_weight =
         one(weight) / (static_cast<remove_complex<ValueType>>(2.0) - weight);
 
+    const auto in_row_ptrs = system_matrix->get_const_row_ptrs();
+    const auto in_col_idxs = system_matrix->get_const_col_idxs();
+    const auto in_values = system_matrix->get_const_values();
+    const auto l_row_ptrs = l_mtx->get_const_row_ptrs();
+    const auto l_col_idxs = l_mtx->get_col_idxs();
+    const auto l_values = l_mtx->get_values();
+    const auto u_row_ptrs = u_mtx->get_const_row_ptrs();
+    const auto u_col_idxs = u_mtx->get_col_idxs();
+    const auto u_values = u_mtx->get_values();
+
     exec->get_queue()->parallel_for(
         sycl_nd_range(grid_dim, block_size), [=](sycl::nd_item<3> item_ct1) {
             factorization::helpers::initialize_l_u(
-                num_rows, system_matrix->get_const_row_ptrs(),
-                system_matrix->get_const_col_idxs(),
-                system_matrix->get_const_values(), l_mtx->get_const_row_ptrs(),
-                l_mtx->get_col_idxs(), l_mtx->get_values(),
-                u_mtx->get_const_row_ptrs(), u_mtx->get_col_idxs(),
-                u_mtx->get_values(),
+                num_rows, in_row_ptrs, in_col_idxs, in_values, l_row_ptrs,
+                l_col_idxs, l_values, u_row_ptrs, u_col_idxs, u_values,
                 factorization::helpers::triangular_mtx_closure(
                     [inv_weight](auto val) { return val * inv_weight; },
                     factorization::helpers::identity{}),

examples/CMakeLists.txt

@@ -101,6 +101,8 @@ if(GINKGO_BUILD_TESTS)
                      "${executor}"
                      WORKING_DIRECTORY
                      "${CMAKE_CURRENT_SOURCE_DIR}/${example}")
+            # Prevent performance issues with high core counts
+            set_property(TEST example_${example}_${executor} PROPERTY ENVIRONMENT OMP_NUM_THREADS=4)
         endforeach()
     endforeach()
 
@@ -115,6 +117,8 @@ if(GINKGO_BUILD_TESTS)
                      "${CMAKE_CURRENT_SOURCE_DIR}/file-config-solver/data/A.mtx"
                      WORKING_DIRECTORY
                      "$<TARGET_FILE_DIR:ginkgo>")
+            # Prevent performance issues with high core counts
+            set_property(TEST example_file-config-solver_${config_name}_${executor} PROPERTY ENVIRONMENT OMP_NUM_THREADS=4)
         endforeach()    
     endforeach()
 

test/factorization/factorization_kernels.cpp

@@ -53,7 +53,7 @@ TEST_F(Factorization, InitializeRowPtrsLSameAsRef)
 }
 
 
-TEST_F(Factorization, InitializeLWithoutSqrtSameAsRef)
+TEST_F(Factorization, InitializeLSameAsRef)
 {
     gko::array<index_type> l_ptrs{ref, mtx->get_size()[0] + 1};
     gko::kernels::reference::factorization::initialize_row_ptrs_l(
@@ -73,6 +73,7 @@ TEST_F(Factorization, InitializeLWithoutSqrtSameAsRef)
         gko::kernels::GKO_DEVICE_NAMESPACE::factorization::initialize_l(
             exec, dmtx.get(), dl_mtx.get(), diag_sqrt);
 
-        GKO_ASSERT_MTX_NEAR(l_mtx, dl_mtx, 0.0);
+        GKO_ASSERT_MTX_NEAR(l_mtx, dl_mtx,
+                            diag_sqrt ? r<value_type>::value : 0.0);
     }
 }

test/preconditioner/sor_kernels.cpp

@@ -46,11 +46,17 @@ class Sor : public CommonTestFixture {
         d_mtx->read(md);
 
         result_l->read(md_l);
-        result_l->scale(gko::initialize<Dense>({0.0}, ref));
+        std::fill_n(result_l->get_col_idxs(),
+                    result_l->get_num_stored_elements(), -1);
+        std::fill_n(result_l->get_values(), result_l->get_num_stored_elements(),
+                    gko::nan<value_type>());
         d_result_l = gko::clone(exec, result_l);
 
         result_u->read(md_u);
-        result_u->scale(gko::initialize<Dense>({0.0}, ref));
+        std::fill_n(result_u->get_col_idxs(),
+                    result_u->get_num_stored_elements(), -1);
+        std::fill_n(result_u->get_values(), result_u->get_num_stored_elements(),
+                    gko::nan<value_type>());
         d_result_u = gko::clone(exec, result_u);
     }
 
@@ -73,6 +79,7 @@ TEST_F(Sor, InitializeWeightedLFactorIsSameAsReference)
     gko::kernels::GKO_DEVICE_NAMESPACE::sor::initialize_weighted_l(
         exec, d_mtx.get(), 1.24, d_result_l.get());
 
+    GKO_ASSERT_MTX_EQ_SPARSITY(result_l, d_result_l);
     GKO_ASSERT_MTX_NEAR(result_l, d_result_l, r<value_type>::value);
 }
 
@@ -84,6 +91,8 @@ TEST_F(Sor, InitializeWeightedLAndUFactorIsSameAsReference)
     gko::kernels::GKO_DEVICE_NAMESPACE::sor::initialize_weighted_l_u(
         exec, d_mtx.get(), 1.24, d_result_l.get(), d_result_u.get());
 
+    GKO_ASSERT_MTX_EQ_SPARSITY(result_l, d_result_l);
+    GKO_ASSERT_MTX_EQ_SPARSITY(result_u, d_result_u);
     GKO_ASSERT_MTX_NEAR(result_l, d_result_l, r<value_type>::value);
     GKO_ASSERT_MTX_NEAR(result_u, d_result_u, r<value_type>::value);
 }