Csric0 threadfence fix (#672) (#360)

* Fix csric0 failures * Added dense diagonally dominant tests to csric0 * Remove hardware tag from yaml file --------- Co-authored-by: jsandham <[email protected]>
ROCm · Nov 29, 2023 · 456a62e · 456a62e
1 parent bb6171c
commit 456a62e
Show file tree

Hide file tree

Showing 3 changed files with 192 additions and 2 deletions.
diff --git a/clients/testings/testing_csric0.cpp b/clients/testings/testing_csric0.cpp
@@ -497,4 +497,181 @@ INSTANTIATE(float);
 INSTANTIATE(double);
 INSTANTIATE(rocsparse_float_complex);
 INSTANTIATE(rocsparse_double_complex);
-void testing_csric0_extra(const Arguments& arg) {}
+
+void testing_csric0_extra(const Arguments& arg)
+{
+    rocsparse_int             M    = arg.M;
+    rocsparse_int             N    = arg.N;
+    rocsparse_analysis_policy apol = arg.apol;
+    rocsparse_solve_policy    spol = arg.spol;
+    rocsparse_index_base      base = arg.baseA;
+
+    // Create rocsparse handle
+    rocsparse_local_handle handle(arg);
+
+    // Create matrix descriptor
+    rocsparse_local_mat_descr descr;
+
+    // Create matrix info
+    rocsparse_local_mat_info info;
+
+    // Set matrix index base
+    CHECK_ROCSPARSE_ERROR(rocsparse_set_mat_index_base(descr, base));
+
+    rocsparse_int nnz = M * N;
+
+    // Allocate host memory for matrix
+    host_vector<rocsparse_int> hcsr_row_ptr(M + 1);
+    host_vector<rocsparse_int> hcsr_col_ind(nnz);
+    host_vector<float>         hcsr_val(nnz);
+
+    // Create dense matrix
+    hcsr_row_ptr[0] = base;
+    for(rocsparse_int i = 0; i < M; i++)
+    {
+        hcsr_row_ptr[i + 1] = hcsr_row_ptr[i] + N;
+    }
+
+    for(rocsparse_int i = 0; i < M; i++)
+    {
+        rocsparse_int start = hcsr_row_ptr[i] - base;
+
+        for(rocsparse_int j = 0; j < N; j++)
+        {
+            hcsr_col_ind[start + j] = j + base;
+            hcsr_val[start + j]     = random_cached_generator_normal<float>();
+            if(i == j)
+            {
+                hcsr_val[start + j] += N + 1;
+            }
+        }
+    }
+
+    host_vector<float> hcsr_val_gold = hcsr_val;
+
+    // Allocate host memory for vectors
+    host_vector<float>         hcsr_val_1(nnz);
+    host_vector<float>         hcsr_val_2(nnz);
+    host_vector<rocsparse_int> h_analysis_pivot_1(1);
+    host_vector<rocsparse_int> h_analysis_pivot_2(1);
+    host_vector<rocsparse_int> h_analysis_pivot_gold(1);
+    host_vector<rocsparse_int> h_solve_pivot_1(1);
+    host_vector<rocsparse_int> h_solve_pivot_2(1);
+    host_vector<rocsparse_int> h_solve_pivot_gold(1);
+
+    // Allocate device memory
+    device_vector<rocsparse_int> dcsr_row_ptr(M + 1);
+    device_vector<rocsparse_int> dcsr_col_ind(nnz);
+    device_vector<float>         dcsr_val_1(nnz);
+    device_vector<float>         dcsr_val_2(nnz);
+    device_vector<rocsparse_int> d_analysis_pivot_2(1);
+    device_vector<rocsparse_int> d_solve_pivot_2(1);
+
+    // Copy data from CPU to device
+    CHECK_HIP_ERROR(hipMemcpy(
+        dcsr_row_ptr, hcsr_row_ptr, sizeof(rocsparse_int) * (M + 1), hipMemcpyHostToDevice));
+    CHECK_HIP_ERROR(
+        hipMemcpy(dcsr_col_ind, hcsr_col_ind, sizeof(rocsparse_int) * nnz, hipMemcpyHostToDevice));
+    CHECK_HIP_ERROR(hipMemcpy(dcsr_val_1, hcsr_val, sizeof(float) * nnz, hipMemcpyHostToDevice));
+
+    // Obtain required buffer size
+    size_t buffer_size;
+    CHECK_ROCSPARSE_ERROR(rocsparse_csric0_buffer_size<float>(
+        handle, M, nnz, descr, dcsr_val_1, dcsr_row_ptr, dcsr_col_ind, info, &buffer_size));
+
+    void* dbuffer;
+    CHECK_HIP_ERROR(rocsparse_hipMalloc(&dbuffer, buffer_size));
+
+    // Copy data from CPU to device
+    CHECK_HIP_ERROR(hipMemcpy(dcsr_val_2, hcsr_val, sizeof(float) * nnz, hipMemcpyHostToDevice));
+
+    // Perform analysis step
+
+    // Pointer mode host
+    CHECK_ROCSPARSE_ERROR(rocsparse_set_pointer_mode(handle, rocsparse_pointer_mode_host));
+    CHECK_ROCSPARSE_ERROR(rocsparse_csric0_analysis<float>(
+        handle, M, nnz, descr, dcsr_val_1, dcsr_row_ptr, dcsr_col_ind, info, apol, spol, dbuffer));
+    {
+        auto st = rocsparse_csric0_zero_pivot(handle, info, h_analysis_pivot_1);
+        EXPECT_ROCSPARSE_STATUS(st,
+                                (h_analysis_pivot_1[0] != -1) ? rocsparse_status_zero_pivot
+                                                              : rocsparse_status_success);
+    }
+
+    // Sync to force updated pivots
+    CHECK_HIP_ERROR(hipDeviceSynchronize());
+
+    // Pointer mode device
+    CHECK_ROCSPARSE_ERROR(rocsparse_set_pointer_mode(handle, rocsparse_pointer_mode_device));
+    CHECK_ROCSPARSE_ERROR(rocsparse_csric0_analysis<float>(
+        handle, M, nnz, descr, dcsr_val_2, dcsr_row_ptr, dcsr_col_ind, info, apol, spol, dbuffer));
+    EXPECT_ROCSPARSE_STATUS(rocsparse_csric0_zero_pivot(handle, info, d_analysis_pivot_2),
+                            (h_analysis_pivot_1[0] != -1) ? rocsparse_status_zero_pivot
+                                                          : rocsparse_status_success);
+
+    // Sync to force updated pivots
+    CHECK_HIP_ERROR(hipDeviceSynchronize());
+
+    // Perform solve step
+
+    // Pointer mode host
+    CHECK_ROCSPARSE_ERROR(rocsparse_set_pointer_mode(handle, rocsparse_pointer_mode_host));
+    CHECK_ROCSPARSE_ERROR(testing::rocsparse_csric0<float>(
+        handle, M, nnz, descr, dcsr_val_1, dcsr_row_ptr, dcsr_col_ind, info, spol, dbuffer));
+    {
+        auto st = rocsparse_csric0_zero_pivot(handle, info, h_solve_pivot_1);
+        EXPECT_ROCSPARSE_STATUS(st,
+                                (h_solve_pivot_1[0] != -1) ? rocsparse_status_zero_pivot
+                                                           : rocsparse_status_success);
+    }
+
+    // Sync to force updated pivots
+    CHECK_HIP_ERROR(hipDeviceSynchronize());
+
+    // Pointer mode device
+    CHECK_ROCSPARSE_ERROR(rocsparse_set_pointer_mode(handle, rocsparse_pointer_mode_device));
+    CHECK_ROCSPARSE_ERROR(testing::rocsparse_csric0<float>(
+        handle, M, nnz, descr, dcsr_val_2, dcsr_row_ptr, dcsr_col_ind, info, spol, dbuffer));
+    EXPECT_ROCSPARSE_STATUS(rocsparse_csric0_zero_pivot(handle, info, d_solve_pivot_2),
+                            (h_solve_pivot_1[0] != -1) ? rocsparse_status_zero_pivot
+                                                       : rocsparse_status_success);
+
+    // Sync to force updated pivots
+    CHECK_HIP_ERROR(hipDeviceSynchronize());
+
+    // Copy output to host
+    CHECK_HIP_ERROR(hipMemcpy(hcsr_val_1, dcsr_val_1, sizeof(float) * nnz, hipMemcpyDeviceToHost));
+    CHECK_HIP_ERROR(hipMemcpy(hcsr_val_2, dcsr_val_2, sizeof(float) * nnz, hipMemcpyDeviceToHost));
+    CHECK_HIP_ERROR(hipMemcpy(
+        h_analysis_pivot_2, d_analysis_pivot_2, sizeof(rocsparse_int), hipMemcpyDeviceToHost));
+    CHECK_HIP_ERROR(
+        hipMemcpy(h_solve_pivot_2, d_solve_pivot_2, sizeof(rocsparse_int), hipMemcpyDeviceToHost));
+
+    // CPU csric0
+    host_csric0<float>(M,
+                       hcsr_row_ptr,
+                       hcsr_col_ind,
+                       hcsr_val_gold,
+                       base,
+                       h_analysis_pivot_gold,
+                       h_solve_pivot_gold);
+
+    // Check pivots
+    h_analysis_pivot_gold.unit_check(h_analysis_pivot_1);
+    h_analysis_pivot_gold.unit_check(h_analysis_pivot_2);
+    h_solve_pivot_gold.unit_check(h_solve_pivot_1);
+    h_solve_pivot_gold.unit_check(h_solve_pivot_2);
+
+    // Check solution vector if no pivot has been found
+    if(h_analysis_pivot_gold[0] == -1 && h_solve_pivot_gold[0] == -1)
+    {
+        hcsr_val_gold.near_check(hcsr_val_1);
+        hcsr_val_gold.near_check(hcsr_val_2);
+    }
+
+    // Clear csric0 meta data
+    CHECK_ROCSPARSE_ERROR(rocsparse_csric0_clear(handle, info));
+
+    // Free buffer
+    CHECK_HIP_ERROR(rocsparse_hipFree(dbuffer));
+}
diff --git a/clients/tests/test_csric0.yaml b/clients/tests/test_csric0.yaml
@@ -41,12 +41,25 @@ Definitions:
     - { M:  37017, N:  37017 }
     - { M: 505194, N: 505194 }
 
+  - &M_N_range_extra
+    - { M: 10, N: 10 }
+    - { M: 235, N: 235 }
+    - { M: 1200, N: 1200 }
+
 Tests:
 - name: csric0_bad_arg
   category: pre_checkin
   function: csric0_bad_arg
   precision: *single_double_precisions_complex_real
 
+- name: csric0_extra
+  category: quick
+  M_N: *M_N_range_extra
+  apol: [rocsparse_analysis_policy_reuse, rocsparse_analysis_policy_force]
+  spol: [rocsparse_solve_policy_auto]
+  baseA: [rocsparse_index_base_zero, rocsparse_index_base_one]
+  function: csric0_extra
+
 - name: csric0
   category: quick
   function: csric0

diff --git a/library/src/precond/csric0_device.h b/library/src/precond/csric0_device.h
@@ -390,6 +390,6 @@ void csric0_binsearch_kernel(rocsparse_int m,
     if(lid == WFSIZE - 1)
     {
         // Last lane writes "we are done" flag
-        __hip_atomic_store(&done[row], 1, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+        __hip_atomic_store(&done[row], 1, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_AGENT);
     }
 }