Add Batch jacobi device kernels #1600

pratikvn · 2024-04-18T22:11:36Z

This PR adds the device kernels for batch Jacobi preconditioner.

MarcelKoch

Overall looks good. Mostly minor comments regarding consistency left.

One thing I've noticed is that some of the functions in the device preconditioner classes explicitly synchronize at the end, while others don't. Maybe it would be best to find a consistent behavior. I think I would prefer to not synchronize in the functions, but instead require synchronization after the function call.

test/preconditioner/batch_jacobi_kernels.cpp

common/cuda_hip/preconditioner/batch_scalar_jacobi.hpp.inc

dpcpp/preconditioner/batch_preconditioners.hpp

common/cuda_hip/preconditioner/batch_block_jacobi.hpp.inc

MarcelKoch · 2024-05-08T11:02:25Z

dpcpp/preconditioner/batch_jacobi_kernels.dp.cpp

+    //    constexpr int subwarp_size =
+    //        gko::kernels::dpcpp::jacobi::get_larger_power(compiled_max_block_size);
+    // TODO: Find a way to allow smaller block_sizes (<16)


this sounds like the whole kernel selection approach is unnecessary here. So maybe remove the template parameter for now and add it later.

dpcpp/preconditioner/batch_jacobi_kernels.hpp.inc

stale

yhmtsai

one major thing left: dealing with d value in invert_dense_block are different between sycl and cuda.

common/cuda_hip/preconditioner/batch_block_jacobi.hpp.inc

common/cuda_hip/preconditioner/batch_jacobi_kernels.hpp.inc

dpcpp/preconditioner/batch_jacobi_kernels.dp.cpp

yhmtsai · 2024-05-13T12:19:54Z

dpcpp/preconditioner/batch_block_jacobi.hpp.inc

+        common_generate_for_all_system_matrix_types(batch_id);
+        item_ct1.barrier(sycl::access::fence_space::local_space);


the barrier should take quite some overhead.
If it is unnecessary, maybe remove it.

yhmtsai · 2024-05-13T12:22:56Z

test/preconditioner/batch_jacobi_kernels.cpp

@@ -257,7 +249,7 @@ TEST_F(BatchJacobi, CanSolveLargeMatrixSizeHpdSystemWithScalarJacobi)
    for (size_t i = 0; i < num_batch_items; i++) {
        auto comp_res_norm = res.host_res_norm->get_const_values()[i] /
                             linear_system.host_rhs_norm->get_const_values()[i];
-        ASSERT_LE(iter_counts->get_const_data()[i], max_iters);
+        EXPECT_LT(iter_counts->get_const_data()[i], max_iters);


If it possbly equal, maybe use EXPECT_LE not EXPECT_LT?

pratikvn · 2024-05-13T13:57:11Z

Will merge this only once #1396 has been merged.

tcojean · 2024-05-22T11:16:00Z

@pratikvn I will check if I can find the issue with the SYCL backend. If I'm not successful, we can just not merge that.

codecov · 2024-05-28T17:56:03Z

Codecov Report

Attention: Patch coverage is 84.90566% with 8 lines in your changes are missing coverage. Please review.

Project coverage is 89.28%. Comparing base (1fc3d7d) to head (d7bc39e).
Report is 13 commits behind head on develop.

❗ Current head d7bc39e differs from pull request most recent head 691a479

Please upload reports for the commit 691a479 to get more accurate results.

Files	Patch %	Lines
reference/preconditioner/batch_scalar_jacobi.hpp	40.00%	6 Missing ⚠️
reference/preconditioner/batch_block_jacobi.hpp	66.66%	1 Missing ⚠️
test/preconditioner/batch_jacobi_kernels.cpp	97.22%	1 Missing ⚠️

Additional details and impacted files

@@             Coverage Diff             @@
##           develop    #1600      +/-   ##
===========================================
- Coverage    89.97%   89.28%   -0.70%     
===========================================
  Files          752      752              
  Lines        60445    60467      +22     
===========================================
- Hits         54387    53985     -402     
- Misses        6058     6482     +424

☔ View full report in Codecov by Sentry.
📢 Have feedback on the report? Share it here.

Co-authored-by: Isha Aggarwal <[email protected]> Co-authored-by: Aditya Kashi <[email protected]>

Co-authored-by: Phuong Nguyen <[email protected]>

Co-authored-by: Marcel Koch <[email protected]>

Co-authored-by: Yu-Hsiang Tsai <[email protected]>

The failure is fairly minimal as it happens only on some integrated GPUs (some Gen 11). Nonetheless, similarly to the DPC++ non-batch Jacobi kernels, this indicates a bigger investigation and fix of these kernels is necessary.

sonarcloud · 2024-05-29T16:47:55Z

Quality Gate passed

Issues
7 New issues
0 Accepted issues

Measures
0 Security Hotspots
75.9% Coverage on New Code
9.6% Duplication on New Code

See analysis details on SonarCloud

pratikvn added the type:batched-functionality This is related to the batched functionality in Ginkgo label Apr 18, 2024

pratikvn self-assigned this Apr 18, 2024

ginkgo-bot added reg:build This is related to the build system. mod:cuda This is related to the CUDA module. type:preconditioner This is related to the preconditioners mod:hip This is related to the HIP module. labels Apr 18, 2024

MarcelKoch added this to the Ginkgo 1.8.0 milestone Apr 19, 2024

pratikvn force-pushed the batch-jacobi-device branch 2 times, most recently from 96c2138 to 5d26c0e Compare April 26, 2024 12:48

pratikvn changed the title ~~WIP: Batch jacobi device kernels~~ Add Batch jacobi device kernels Apr 26, 2024

pratikvn added the 1:ST:ready-for-review This PR is ready for review label Apr 26, 2024

pratikvn requested review from a team April 26, 2024 12:50

pratikvn force-pushed the batch-jacobi-device branch 3 times, most recently from 810bc89 to c5c918e Compare April 30, 2024 08:26

pratikvn force-pushed the batch-prec-jacobi branch from a2e3afa to 4a6e2ef Compare May 2, 2024 21:43

MarcelKoch requested review from MarcelKoch and yhmtsai May 3, 2024 09:26

pratikvn force-pushed the batch-prec-jacobi branch 2 times, most recently from 32de29b to fcff54f Compare May 6, 2024 22:59

pratikvn force-pushed the batch-jacobi-device branch from c5c918e to f966541 Compare May 6, 2024 23:00

MarcelKoch previously requested changes May 8, 2024

View reviewed changes

pratikvn force-pushed the batch-prec-jacobi branch from 5504e62 to 83a9280 Compare May 9, 2024 13:27

Base automatically changed from batch-prec-jacobi to develop May 9, 2024 16:53

pratikvn force-pushed the batch-jacobi-device branch from 05eda68 to 2911666 Compare May 9, 2024 20:32

pratikvn requested a review from MarcelKoch May 9, 2024 20:42

pratikvn force-pushed the batch-jacobi-device branch from dab3875 to 92c6ba0 Compare May 9, 2024 20:44

pratikvn added the 1:ST:no-changelog-entry Skip the wiki check for changelog update label May 9, 2024

yhmtsai requested changes May 13, 2024

View reviewed changes

pratikvn requested a review from yhmtsai May 13, 2024 13:35

yhmtsai approved these changes May 13, 2024

View reviewed changes

pratikvn added 1:ST:ready-to-merge This PR is ready to merge. and removed 1:ST:ready-for-review This PR is ready for review labels May 13, 2024

pratikvn force-pushed the batch-jacobi-device branch 3 times, most recently from 973c77b to 37e71d8 Compare May 16, 2024 21:20

tcojean force-pushed the batch-jacobi-device branch from 37e71d8 to d7bc39e Compare May 28, 2024 06:48

pratikvn and others added 14 commits May 29, 2024 08:58

add cuda/hip kernels and tests

6c38777

Co-authored-by: Isha Aggarwal <[email protected]> Co-authored-by: Aditya Kashi <[email protected]>

test updates

19f36e0

Fix tests

ef4726c

Add dpcpp kernels

37d2cd8

Co-authored-by: Phuong Nguyen <[email protected]>

dpcpp kernel fixes

fbb8ec5

test fixes

b2a41ac

review updates

21092b2

Co-authored-by: Marcel Koch <[email protected]>

fix typos

67d0aa3

update using decls, and fix dyn work func

dddcec1

comment format

e22a6b7

review updates

4760455

Co-authored-by: Marcel Koch <[email protected]>

review updates

1c3d1a0

Co-authored-by: Yu-Hsiang Tsai <[email protected]>

ensure block jac diag is not zero

202f227

batch jacobi, dpcpp: disable failing tests

691a479

The failure is fairly minimal as it happens only on some integrated GPUs (some Gen 11). Nonetheless, similarly to the DPC++ non-batch Jacobi kernels, this indicates a bigger investigation and fix of these kernels is necessary.

tcojean force-pushed the batch-jacobi-device branch from d7bc39e to 691a479 Compare May 29, 2024 06:58

tcojean merged commit 1547509 into develop May 29, 2024
11 of 15 checks passed

tcojean deleted the batch-jacobi-device branch May 29, 2024 12:05

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add Batch jacobi device kernels #1600

Add Batch jacobi device kernels #1600

pratikvn commented Apr 18, 2024

MarcelKoch left a comment

MarcelKoch May 8, 2024

yhmtsai left a comment

yhmtsai May 13, 2024

yhmtsai May 13, 2024

pratikvn commented May 13, 2024

tcojean commented May 22, 2024 •

edited

Loading

codecov bot commented May 28, 2024 •

edited

Loading

sonarcloud bot commented May 29, 2024

		common_generate_for_all_system_matrix_types(batch_id);
		item_ct1.barrier(sycl::access::fence_space::local_space);

Add Batch jacobi device kernels #1600

Add Batch jacobi device kernels #1600

Conversation

pratikvn commented Apr 18, 2024

MarcelKoch left a comment

Choose a reason for hiding this comment

MarcelKoch May 8, 2024

Choose a reason for hiding this comment

yhmtsai left a comment

Choose a reason for hiding this comment

yhmtsai May 13, 2024

Choose a reason for hiding this comment

yhmtsai May 13, 2024

Choose a reason for hiding this comment

pratikvn commented May 13, 2024

tcojean commented May 22, 2024 • edited Loading

codecov bot commented May 28, 2024 • edited Loading

Codecov Report

sonarcloud bot commented May 29, 2024

Quality Gate passed

tcojean commented May 22, 2024 •

edited

Loading

codecov bot commented May 28, 2024 •

edited

Loading