Add NVHPC CI and fix issues #1331

tcojean · 2023-05-04T16:27:47Z

This is no longer a WIP, but issues remain:

~~Investigate matrix_generator_test likely miscompilation~~ I just disabled the offending tests, if somebody wants to investigate, check why Release builds before disable miscompiling matrix_generator tests fail.
Investigate jacobi_kernels adaptive/reduced precision issues

cuda/test/solver/upper_trs_kernels.cpp

include/ginkgo/core/base/std_extensions.hpp

cuda/test/solver/lower_trs_kernels.cpp

yhmtsai · 2023-05-14T14:58:36Z

omp/solver/idr_kernels.cpp

-#pragma omp parallel for reduction(+ : norm)
-        for (size_type j = 0; j < num_cols; j++) {
-            norm += squared_norm(subspace_vectors->at(row, j));
-        }
+        run_kernel_reduction(
+            exec,
+            [](auto col, auto row, auto subspace_vectors) {
+                return squared_norm(subspace_vectors(row, col));
+            },
+            GKO_KERNEL_REDUCE_SUM(remove_complex<ValueType>), &norm, num_cols,
+            static_cast<int64>(row), subspace_vectors);


Is it for fix or optimization?

It's a first step towards unifying it, but also a workaround for NVHPC's limited pragma omp reduction feature set

Do you have any specific feature name for the unsupported one and the corresponding error message?

basically pragma omp (parallel) for reduction is supported, but omp declare reduction is not

I see. because omp run_kernel_reduction does not use omp reduction (declare for complex), it works via run_kernel_reduction

.gitlab-ci.yml

.gitlab/image.yml

yhmtsai

others LGTM. Some of CI tests are not passed, so I do not put my approval yet

yhmtsai · 2023-05-16T10:24:27Z

omp/solver/idr_kernels.cpp

-#pragma omp parallel for reduction(+ : norm)
-        for (size_type j = 0; j < num_cols; j++) {
-            norm += squared_norm(subspace_vectors->at(row, j));
-        }
+        run_kernel_reduction(
+            exec,
+            [](auto col, auto row, auto subspace_vectors) {
+                return squared_norm(subspace_vectors(row, col));
+            },
+            GKO_KERNEL_REDUCE_SUM(remove_complex<ValueType>), &norm, num_cols,
+            static_cast<int64>(row), subspace_vectors);


Do you have any specific feature name for the unsupported one and the corresponding error message?

CUDA >= 11.7 seems to use new kernels for all devices

tcojean · 2023-06-05T16:49:18Z

format!

tcojean · 2023-06-05T16:56:54Z

This should now be ""working"" including for Jacobi. I am also trying the new containers at the same time with the fixed environment.

yhmtsai

LGTM. only have comments on testing version - should keep at least one job for the minimal supported version like cuda 9.2 and dpcpp 2022.1

.gitlab-ci.yml

yhmtsai · 2023-06-06T18:39:47Z

omp/solver/idr_kernels.cpp

-#pragma omp parallel for reduction(+ : norm)
-        for (size_type j = 0; j < num_cols; j++) {
-            norm += squared_norm(subspace_vectors->at(row, j));
-        }
+        run_kernel_reduction(
+            exec,
+            [](auto col, auto row, auto subspace_vectors) {
+                return squared_norm(subspace_vectors(row, col));
+            },
+            GKO_KERNEL_REDUCE_SUM(remove_complex<ValueType>), &norm, num_cols,
+            static_cast<int64>(row), subspace_vectors);


I see. because omp run_kernel_reduction does not use omp reduction (declare for complex), it works via run_kernel_reduction

upsj

LGTM (though somebody else should review my changes)

omp/CMakeLists.txt

codecov · 2023-06-07T08:54:00Z

Codecov Report

Patch coverage: 100.00% and project coverage change: -0.39 ⚠️

Comparison is base (d94433b) 91.19% compared to head (75b8c52) 90.81%.

Additional details and impacted files

@@             Coverage Diff             @@
##           develop    #1331      +/-   ##
===========================================
- Coverage    91.19%   90.81%   -0.39%     
===========================================
  Files          598      598              
  Lines        50628    50666      +38     
===========================================
- Hits         46172    46011     -161     
- Misses        4456     4655     +199

Impacted Files	Coverage Δ
include/ginkgo/core/base/matrix_assembly_data.hpp	`100.00% <ø> (ø)`
include/ginkgo/core/base/polymorphic_object.hpp	`89.90% <ø> (ø)`
include/ginkgo/core/log/logger.hpp	`93.93% <ø> (ø)`
reference/test/solver/cgs_kernels.cpp	`100.00% <ø> (ø)`
reference/test/solver/gcr_kernels.cpp	`99.19% <ø> (-0.01%)`	⬇️
core/test/utils/matrix_generator_test.cpp	`97.93% <100.00%> (-2.07%)`	⬇️
omp/distributed/partition_kernels.cpp	`100.00% <100.00%> (ø)`
omp/solver/cb_gmres_kernels.cpp	`92.30% <100.00%> (+1.16%)`	⬆️
omp/solver/idr_kernels.cpp	`75.00% <100.00%> (ø)`

... and 5 files with indirect coverage changes

☔ View full report in Codecov by Sentry.
📢 Do you have feedback about the report comment? Let us know in this issue.

Co-authored-by: Terry Cojean <tcojean@users.noreply.github.com>

sonarcloud · 2023-06-07T22:20:13Z

Kudos, SonarCloud Quality Gate passed!

0 Bugs
0 Vulnerabilities
0 Security Hotspots
9 Code Smells

8.7% Coverage
0.0% Duplication

Release 1.6.0 of Ginkgo. The Ginkgo team is proud to announce the new Ginkgo minor release 1.6.0. This release brings new features such as: - Several building blocks for GPU-resident sparse direct solvers like symbolic and numerical LU and Cholesky factorization, ..., - A distributed Schwarz preconditioner, - New FGMRES and GCR solvers, - Distributed benchmarks for the SpMV operation, solvers, ... - Support for non-default streams in the CUDA and HIP backends, - Mixed precision support for the CSR SpMV, - A new profiling logger which integrates with NVTX, ROCTX, TAU and VTune to provide internal Ginkgo knowledge to most HPC profilers! and much more. If you face an issue, please first check our [known issues page](https://github.com/ginkgo-project/ginkgo/wiki/Known-Issues) and the [open issues list](https://github.com/ginkgo-project/ginkgo/issues) and if you do not find a solution, feel free to [open a new issue](https://github.com/ginkgo-project/ginkgo/issues/new/choose) or ask a question using the [github discussions](https://github.com/ginkgo-project/ginkgo/discussions). Supported systems and requirements: + For all platforms, CMake 3.13+ + C++14 compliant compiler + Linux and macOS + GCC: 5.5+ + clang: 3.9+ + Intel compiler: 2018+ + Apple Clang: 14.0 is tested. Earlier versions might also work. + NVHPC: 22.7+ + Cray Compiler: 14.0.1+ + CUDA module: CUDA 9.2+ or NVHPC 22.7+ + HIP module: ROCm 4.5+ + DPC++ module: Intel OneAPI 2021.3+ with oneMKL and oneDPL. Set the CXX compiler to `dpcpp`. + Windows + MinGW: GCC 5.5+ + Microsoft Visual Studio: VS 2019+ + CUDA module: CUDA 9.2+, Microsoft Visual Studio + OpenMP module: MinGW. ### Version Support Changes + ROCm 4.0+ -> 4.5+ after [#1303](#1303) + Removed Cygwin pipeline and support [#1283](#1283) ### Interface Changes + Due to internal changes, `ConcreteExecutor::run` will now always throw if the corresponding module for the `ConcreteExecutor` is not build [#1234](#1234) + The constructor of `experimental::distributed::Vector` was changed to only accept local vectors as `std::unique_ptr` [#1284](#1284) + The default parameters for the `solver::MultiGrid` were improved. In particular, the smoother defaults to one iteration of `Ir` with `Jacobi` preconditioner, and the coarse grid solver uses the new direct solver with LU factorization. [#1291](#1291) [#1327](#1327) + The `iteration_complete` event gained a more expressive overload with additional parameters, the old overloads were deprecated. [#1288](#1288) [#1327](#1327) ### Deprecations + Deprecated less expressive `iteration_complete` event. Users are advised to now implement the function `void iteration_complete(const LinOp* solver, const LinOp* b, const LinOp* x, const size_type& it, const LinOp* r, const LinOp* tau, const LinOp* implicit_tau_sq, const array<stopping_status>* status, bool stopped)` [#1288](#1288) ### Added Features + A distributed Schwarz preconditioner. [#1248](#1248) + A GCR solver [#1239](#1239) + Flexible Gmres solver [#1244](#1244) + Enable Gmres solver for distributed matrices and vectors [#1201](#1201) + An example that uses Kokkos to assemble the system matrix [#1216](#1216) + A symbolic LU factorization allowing the `gko::experimental::factorization::Lu` and `gko::experimental::solver::Direct` classes to be used for matrices with non-symmetric sparsity pattern [#1210](#1210) + A numerical Cholesky factorization [#1215](#1215) + Symbolic factorizations in host-side operations are now wrapped in a host-side `Operation` to make their execution visible to loggers. This means that profiling loggers and benchmarks are no longer missing a separate entry for their runtime [#1232](#1232) + Symbolic factorization benchmark [#1302](#1302) + The `ProfilerHook` logger allows annotating the Ginkgo execution (apply, operations, ...) for profiling frameworks like NVTX, ROCTX and TAU. [#1055](#1055) + `ProfilerHook::created_(nested_)summary` allows the generation of a lightweight runtime profile over all Ginkgo functions written to a user-defined stream [#1270](#1270) for both host and device timing functionality [#1313](#1313) + It is now possible to enable host buffers for MPI communications at runtime even if the compile option `GINKGO_FORCE_GPU_AWARE_MPI` is set. [#1228](#1228) + A stencil matrices generator (5-pt, 7-pt, 9-pt, and 27-pt) for benchmarks [#1204](#1204) + Distributed benchmarks (multi-vector blas, SpMV, solver) [#1204](#1204) + Benchmarks for CSR sorting and lookup [#1219](#1219) + A timer for MPI benchmarks that reports the longest time [#1217](#1217) + A `timer_method=min|max|average|median` flag for benchmark timing summary [#1294](#1294) + Support for non-default streams in CUDA and HIP executors [#1236](#1236) + METIS integration for nested dissection reordering [#1296](#1296) + SuiteSparse AMD integration for fillin-reducing reordering [#1328](#1328) + Csr mixed-precision SpMV support [#1319](#1319) + A `with_loggers` function for all `Factory` parameters [#1337](#1337) ### Improvements + Improve naming of kernel operations for loggers [#1277](#1277) + Annotate solver iterations in `ProfilerHook` [#1290](#1290) + Allow using the profiler hooks and inline input strings in benchmarks [#1342](#1342) + Allow passing smart pointers in place of raw pointers to most matrix functions. This means that things like `vec->compute_norm2(x.get())` or `vec->compute_norm2(lend(x))` can be simplified to `vec->compute_norm2(x)` [#1279](#1279) [#1261](#1261) + Catch overflows in prefix sum operations, which makes Ginkgo's operations much less likely to crash. This also improves the performance of the prefix sum kernel [#1303](#1303) + Make the installed GinkgoConfig.cmake file relocatable and follow more best practices [#1325](#1325) ### Fixes + Fix OpenMPI version check [#1200](#1200) + Fix the mpi cxx type binding by c binding [#1306](#1306) + Fix runtime failures for one-sided MPI wrapper functions observed on some OpenMPI versions [#1249](#1249) + Disable thread pinning with GPU executors due to poor performance [#1230](#1230) + Fix hwloc version detection [#1266](#1266) + Fix PAPI detection in non-implicit include directories [#1268](#1268) + Fix PAPI support for newer PAPI versions: [#1321](#1321) + Fix pkg-config file generation for library paths outside prefix [#1271](#1271) + Fix various build failures with ROCm 5.4, CUDA 12, and OneAPI 6 [#1214](#1214), [#1235](#1235), [#1251](#1251) + Fix incorrect read for skew-symmetric MatrixMarket files with explicit diagonal entries [#1272](#1272) + Fix handling of missing diagonal entries in symbolic factorizations [#1263](#1263) + Fix segmentation fault in benchmark matrix construction [#1299](#1299) + Fix the stencil matrix creation for benchmarking [#1305](#1305) + Fix the additional residual check in IR [#1307](#1307) + Fix the cuSPARSE CSR SpMM issue on single strided vector when cuda >= 11.6 [#1322](#1322) [#1331](#1331) + Fix Isai generation for large sparsity powers [#1327](#1327) + Fix Ginkgo compilation and test with NVHPC >= 22.7 [#1331](#1331) + Fix Ginkgo compilation of 32 bit binaries with MSVC [#1349](#1349)

ginkgo-bot added reg:ci-cd This is related to the continuous integration system. mod:core This is related to the core module. labels May 4, 2023

tcojean force-pushed the fix_nvhpc branch 2 times, most recently from 2be1dbf to 18c9470 Compare May 5, 2023 09:12

pratikvn added this to the Release 1.6.0 milestone May 8, 2023

upsj force-pushed the fix_nvhpc branch from 18c9470 to 68ae859 Compare May 12, 2023 21:35

upsj changed the title ~~WIP: Add NVHPC CI and fix issues~~ Add NVHPC CI and fix issues May 12, 2023

upsj added the 1:ST:ready-for-review This PR is ready for review label May 12, 2023

upsj requested a review from a team May 12, 2023 21:38

pratikvn reviewed May 13, 2023

View reviewed changes

cuda/test/solver/upper_trs_kernels.cpp Show resolved Hide resolved

include/ginkgo/core/base/std_extensions.hpp Outdated Show resolved Hide resolved

upsj added the 1:ST:run-full-test label May 13, 2023

tcojean commented May 14, 2023

View reviewed changes

cuda/test/solver/lower_trs_kernels.cpp Show resolved Hide resolved

yhmtsai reviewed May 14, 2023

View reviewed changes

yhmtsai reviewed May 16, 2023

View reviewed changes

.gitlab/image.yml Show resolved Hide resolved

yhmtsai reviewed May 16, 2023

View reviewed changes

tcojean and others added 11 commits June 5, 2023 17:01

Add NVHPC jobs. Disable gflags warnings.

2d95059

Fix nvhpc compilation

2fb2e8d

Adapt to std::uncaught_exception() deprecation

f85b5ff

Make nvhpc math abide IEEE standard

278151e

fix nvc++ OpenMP support for IDR and CB-GMRES

88609f5

work around cuSPARSE bug

26f548e

fall back to Csr SpMV for more cuSPARSE bugs

675d3a8

CUDA >= 11.7 seems to use new kernels for all devices

increase test tolerances for nvc++

1307f04

disable miscompiling matrix_generator tests

19ea58f

Fix OpenMP partition kernels

f0d8f5b

Reduce warning amount with nvhpc

fc3786b

tcojean force-pushed the fix_nvhpc branch from bd3b858 to ac27061 Compare June 5, 2023 16:47

Workaround NVHPC jacobi mantissa segmentation bug

9766541

tcojean force-pushed the fix_nvhpc branch from 7652d96 to 009cc4c Compare June 6, 2023 15:26

yhmtsai requested changes Jun 6, 2023

View reviewed changes

upsj approved these changes Jun 7, 2023

View reviewed changes

omp/CMakeLists.txt Show resolved Hide resolved

ginkgo-bot and others added 3 commits June 7, 2023 11:18

Format files

09749bb

Co-authored-by: Terry Cojean <tcojean@users.noreply.github.com>

Put back CUDA 9.2 job for now.

9cb4b34

Try to use oneAPI 2022.1 for one job

75b8c52

tcojean force-pushed the fix_nvhpc branch from 009cc4c to 75b8c52 Compare June 7, 2023 11:24

yhmtsai approved these changes Jun 7, 2023

View reviewed changes

tcojean merged commit 2db15ed into develop Jun 8, 2023

tcojean deleted the fix_nvhpc branch June 8, 2023 16:32

This was referenced Jun 8, 2023

Couple Ginkgo tests are failing on NERSC Perlmutter #1292

Closed

Problems compiling with the NVHPC suite #941

Closed

This was referenced Dec 5, 2023

ginkgo 1.7.0: Compilation with Visual Studio 2022 and C++20 fails 'uncaught_exception': is not a member of 'std' #1495

Closed

fix MSVC __cplusplus by _MSVC_LANG #1496

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add NVHPC CI and fix issues #1331

Add NVHPC CI and fix issues #1331

tcojean commented May 4, 2023 •

edited

Loading

yhmtsai May 14, 2023

upsj May 16, 2023

yhmtsai May 16, 2023

upsj May 16, 2023

yhmtsai Jun 6, 2023

yhmtsai left a comment

yhmtsai May 16, 2023

tcojean commented Jun 5, 2023

tcojean commented Jun 5, 2023

yhmtsai left a comment

yhmtsai Jun 6, 2023

upsj left a comment

codecov bot commented Jun 7, 2023 •

edited

Loading

sonarcloud bot commented Jun 7, 2023

Add NVHPC CI and fix issues #1331

Add NVHPC CI and fix issues #1331

Conversation

tcojean commented May 4, 2023 • edited Loading

yhmtsai May 14, 2023

Choose a reason for hiding this comment

upsj May 16, 2023

Choose a reason for hiding this comment

yhmtsai May 16, 2023

Choose a reason for hiding this comment

upsj May 16, 2023

Choose a reason for hiding this comment

yhmtsai Jun 6, 2023

Choose a reason for hiding this comment

yhmtsai left a comment

Choose a reason for hiding this comment

yhmtsai May 16, 2023

Choose a reason for hiding this comment

tcojean commented Jun 5, 2023

tcojean commented Jun 5, 2023

yhmtsai left a comment

Choose a reason for hiding this comment

yhmtsai Jun 6, 2023

Choose a reason for hiding this comment

upsj left a comment

Choose a reason for hiding this comment

codecov bot commented Jun 7, 2023 • edited Loading

Codecov Report

sonarcloud bot commented Jun 7, 2023

tcojean commented May 4, 2023 •

edited

Loading

codecov bot commented Jun 7, 2023 •

edited

Loading