Simple kernel reduction #833

upsj · 2021-07-13T17:26:12Z

This adds reduction kernels to the simple kernel setup. They use the same setup as the normal kernels, only the inner lambda returns a value over which we will reduce, and a description of the reduction operation, namely identity value (0 for summation), reduction operator (+ for summation) and finalize function (identity for summation, e.g. sqrt for norm2). I will probably add simplifying overloads for summation reduction, since that is the most common operation (next to sqrt of sum and maximum)

TODO:

DPCPP kernels
Tune GPU heuristics
Rowwise/columnwise operations

sonarcloud · 2021-08-13T06:19:34Z

SonarCloud Quality Gate failed.

0 Bugs
0 Vulnerabilities
0 Security Hotspots
37 Code Smells

75.5% Coverage
23.4% Duplication

codecov · 2021-08-13T06:22:27Z

Codecov Report

Merging #833 (dfee616) into develop (b218bb0) will increase coverage by 0.01%.
The diff coverage is 96.39%.

@@             Coverage Diff             @@
##           develop     #833      +/-   ##
===========================================
+ Coverage    94.72%   94.73%   +0.01%     
===========================================
  Files          430      431       +1     
  Lines        35503    35668     +165     
===========================================
+ Hits         33631    33791     +160     
- Misses        1872     1877       +5

Impacted Files	Coverage Δ
omp/matrix/dense_kernels.cpp	`97.41% <ø> (-0.21%)`	⬇️
omp/base/kernel_launch.hpp	`84.61% <88.23%> (-4.97%)`	⬇️
omp/base/kernel_launch_reduction.hpp	`96.62% <96.62%> (ø)`
omp/test/base/kernel_launch.cpp	`95.50% <96.90%> (+1.67%)`	⬆️
common/unified/base/kernel_launch.hpp	`100.00% <100.00%> (ø)`
common/unified/base/kernel_launch_solver.hpp	`90.90% <100.00%> (ø)`
common/unified/matrix/dense_kernels.cpp	`100.00% <100.00%> (ø)`
omp/base/kernel_launch_solver.hpp	`100.00% <100.00%> (ø)`

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update b218bb0...dfee616. Read the comment docs.

yhmtsai

Do not finish the review yet.
cuda/base/kernel_launch_reduction.cuh the kernels seems to be the same as hip/base/kernel_launch_reduction.cuh.
should the init be applied to all elements?
It can not work for summation, right?

dpcpp/base/kernel_launch_solver.dp.hpp

core/test/utils/assertions.hpp

cuda/base/kernel_launch_reduction.cuh

yhmtsai · 2021-08-26T12:39:07Z

cuda/base/kernel_launch_reduction.cuh

+ }
+ partial = reduce(subwarp, partial, op);
+ if (subwarp.thread_rank() == 0) {
+ result[(row + col_block * rows) * result_stride] = finalize(partial);


Is it good for the next level reduction?
Is row * result_stride + col_block allowed, if the result is the internal storage?

for the partial reduction, I set result_stride to 1, is that what you mean?

I thought it is row major so row * cols + col_block?

the final output is a column vector, one entry per row. I'm not sure it really matters whether we store the intermediate results grouped by col_block or by row, I just chose the one that places "near subwarps" into "near memory".

cuda/base/kernel_launch_reduction.cuh

yhmtsai · 2021-08-26T12:49:51Z

cuda/base/kernel_launch_reduction.cuh

+ if (col < cols) {
+ for (auto row = subwarp_id; row < rows; row += subwarp_num) {
+ partial = op(partial, fn(row, col, args...));
+ }
+ }
+ // accumulate between all subwarps in the warp
+#pragma unroll
+ for (unsigned i = subwarp_size; i < warp_size; i *= 2) {
+ partial = op(partial, warp.shfl_xor(partial, i));
+ } // store the result to shared memory


one element of subwarp for one col and use more than 1 subwarp to for loop.
the #subwarp is the parallelization of accumulation and #elements is #cols, right?

yes, the subwarp size is the number of columns rounded up to a power of two, this sums up neighboring subwarps until we have a single result for the whole warp. That works transparently for anything between 1 - 32 columns

cuda/test/base/kernel_launch.cu

thoasm

Only minor comments, LGTM!

common/unified/base/kernel_launch.hpp

core/test/utils/assertions.hpp

cuda/base/kernel_launch_reduction.cuh

yhmtsai

How about to use base or ref not init?
it and the init in #831 will be confusing to me if both they use init.

yhmtsai · 2021-10-04T12:35:28Z

cuda/test/base/kernel_launch.cu

+ [] GKO_KERNEL(auto i, auto a) {
+ static_assert(is_same<decltype(i), int64>::value, "index");
+ static_assert(is_same<decltype(a), int64 *>::value, "value");


Suggested change

[] GKO_KERNEL(auto i, auto a) {

static_assert(is_same<decltype(i), int64>::value, "index");

static_assert(is_same<decltype(a), int64 *>::value, "value");

[] GKO_KERNEL(int64 i, int64* a) {

Could it be this kind?
the first part might have potential conversion, but the second should be detectable?

How could the first one involve conversions? The parameter types are exactly what is being passed in (minus references), while the second case, you could pass in other types.
I want to make sure here that the parameters are exactly the type I am expecting, not something convertible to it.

Yes, I mean the first element might have implicit conversion in the second case

dpcpp/base/kernel_launch_reduction.dp.hpp

yhmtsai · 2021-10-04T12:44:16Z

dpcpp/base/kernel_launch_reduction.dp.hpp

+
+ cgh.parallel_for(
+ range, [=
+ ](sycl::nd_item<3> idx) [[intel::reqd_sub_group_size(sg_size)]] {


sad :(
By the way, is [=] formatted?

yes, clang-format doesn't seem to play nice with attributes right now. With SYCL 2020, maybe we can omit the intel::, but I haven't tried that yet.

yhmtsai · 2021-10-04T13:24:53Z

omp/base/kernel_launch.hpp

 const auto rounded_cols = cols / block_size * block_size;
 GKO_ASSERT(rounded_cols + remainder_cols == cols);
+ if (rounded_cols == 0 || cols == block_size) {
+ // we group all sizes <= block_size here and unroll explicitly


Is this part helpful for the size <= block_size if comparing it against the remainder part in the following?

Yes, the remainder part is just necessary for distinguishing between size == block_size and size < block_size, since remainder_cols only tells us the correct unroll size for size < block_size.

yhmtsai · 2021-10-04T14:15:14Z

omp/base/kernel_launch_reduction.hpp

+
+ auto local_partial = init;
+ if (rounded_cols == 0 || cols == block_size) {
+ // we group all sizes <= block_size here and unroll explicitly


dpcpp/base/kernel_launch_reduction.dp.hpp

yhmtsai · 2021-10-04T14:29:44Z

dpcpp/base/kernel_launch_reduction.dp.hpp

+#pragma unroll
+ for (int i = 1; i < ssg_size; i *= 2) {
+ partial = op(partial, subgroup.shfl_xor(partial, i));
+ }
+ if (col_block < col_blocks && ssg_rank == 0) {
+ result[(row + col_block * rows) * result_stride] =
+ finalize(partial);
+ }
+ });
+ });


I am also thinking add this part into cooperative group subreduce.

yhmtsai · 2021-10-04T15:01:11Z

cuda/base/kernel_launch_reduction.cuh

+ const auto row = idx % rows;
+ const auto col_block = idx / rows;


Does it use column-major?

I'm not sure if that term fits well here - we use one subwarp per row, so the reads are usually coalescing if the input is ready from a row-major matrix.

ah, I see. It is the subwarp id, not the thread id.

yhmtsai · 2021-10-04T15:02:03Z

cuda/base/kernel_launch_reduction.cuh

+ }
+ partial = reduce(subwarp, partial, op);
+ if (subwarp.thread_rank() == 0) {
+ result[(row + col_block * rows) * result_stride] = finalize(partial);


I thought it is row major so row * cols + col_block?

upsj · 2021-10-04T15:26:49Z

@yhmtsai I think base and ref would be a better description for what #831 does. Here, we have the neutral or identity element of the monoid we are reducing over. I went with init initially since that's what std::reduce and std::accumulate use, but there they use slightly different meanings as well, so I am not sure. How does identity sound?

yhmtsai · 2021-10-04T19:15:16Z

std::reduce and std::accumulate use init only once, so I think it makes sense.
Yes, identity sounds good and it fits the mathematics sense

icpc has issues with some combination of pragma omp parallel for and lambdas called inside the loop: internal error: assertion failed: find_assoc_pragma: pragma not found Putting the entire scope into an immediately evaluated lambda helps

* remove unnecessary shmem init * add test comments

Co-authored-by: Tobias Ribizel <[email protected]>

sonarcloud · 2021-10-13T08:32:03Z

SonarCloud Quality Gate failed.

0 Bugs
0 Vulnerabilities
0 Security Hotspots
37 Code Smells

75.5% Coverage
22.9% Duplication

Advertise release 1.5.0 and last changes + Add changelog, + Update third party libraries + A small fix to a CMake file See PR: #1195 The Ginkgo team is proud to announce the new Ginkgo minor release 1.5.0. This release brings many important new features such as: - MPI-based multi-node support for all matrix formats and most solvers; - full DPC++/SYCL support, - functionality and interface for GPU-resident sparse direct solvers, - an interface for wrapping solvers with scaling and reordering applied, - a new algebraic Multigrid solver/preconditioner, - improved mixed-precision support, - support for device matrix assembly, 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 LLVM: 8.0+ + NVHPC: 22.7+ + Cray Compiler: 14.0.1+ + CUDA module: CUDA 9.2+ or NVHPC 22.7+ + HIP module: ROCm 4.0+ + DPC++ module: Intel OneAPI 2021.3 with oneMKL and oneDPL. Set the CXX compiler to `dpcpp`. + Windows + MinGW and Cygwin: GCC 5.5+ + Microsoft Visual Studio: VS 2019 + CUDA module: CUDA 9.2+, Microsoft Visual Studio + OpenMP module: MinGW or Cygwin. Algorithm and important feature additions: + Add MPI-based multi-node for all matrix formats and solvers (except GMRES and IDR). ([#676](#676), [#908](#908), [#909](#909), [#932](#932), [#951](#951), [#961](#961), [#971](#971), [#976](#976), [#985](#985), [#1007](#1007), [#1030](#1030), [#1054](#1054), [#1100](#1100), [#1148](#1148)) + Porting the remaining algorithms (preconditioners like ISAI, Jacobi, Multigrid, ParILU(T) and ParIC(T)) to DPC++/SYCL, update to SYCL 2020, and improve support and performance ([#896](#896), [#924](#924), [#928](#928), [#929](#929), [#933](#933), [#943](#943), [#960](#960), [#1057](#1057), [#1110](#1110), [#1142](#1142)) + Add a Sparse Direct interface supporting GPU-resident numerical LU factorization, symbolic Cholesky factorization, improved triangular solvers, and more ([#957](#957), [#1058](#1058), [#1072](#1072), [#1082](#1082)) + Add a ScaleReordered interface that can wrap solvers and automatically apply reorderings and scalings ([#1059](#1059)) + Add a Multigrid solver and improve the aggregation based PGM coarsening scheme ([#542](#542), [#913](#913), [#980](#980), [#982](#982), [#986](#986)) + Add infrastructure for unified, lambda-based, backend agnostic, kernels and utilize it for some simple kernels ([#833](#833), [#910](#910), [#926](#926)) + Merge different CUDA, HIP, DPC++ and OpenMP tests under a common interface ([#904](#904), [#973](#973), [#1044](#1044), [#1117](#1117)) + Add a device_matrix_data type for device-side matrix assembly ([#886](#886), [#963](#963), [#965](#965)) + Add support for mixed real/complex BLAS operations ([#864](#864)) + Add a FFT LinOp for all but DPC++/SYCL ([#701](#701)) + Add FBCSR support for NVIDIA and AMD GPUs and CPUs with OpenMP ([#775](#775)) + Add CSR scaling ([#848](#848)) + Add array::const_view and equivalent to create constant matrices from non-const data ([#890](#890)) + Add a RowGatherer LinOp supporting mixed precision to gather dense matrix rows ([#901](#901)) + Add mixed precision SparsityCsr SpMV support ([#970](#970)) + Allow creating CSR submatrix including from (possibly discontinuous) index sets ([#885](#885), [#964](#964)) + Add a scaled identity addition (M <- aI + bM) feature interface and impls for Csr and Dense ([#942](#942)) Deprecations and important changes: + Deprecate AmgxPgm in favor of the new Pgm name. ([#1149](#1149)). + Deprecate specialized residual norm classes in favor of a common `ResidualNorm` class ([#1101](#1101)) + Deprecate CamelCase non-polymorphic types in favor of snake_case versions (like array, machine_topology, uninitialized_array, index_set) ([#1031](#1031), [#1052](#1052)) + Bug fix: restrict gko::share to rvalue references (*possible interface break*) ([#1020](#1020)) + Bug fix: when using cuSPARSE's triangular solvers, specifying the factory parameter `num_rhs` is now required when solving for more than one right-hand side, otherwise an exception is thrown ([#1184](#1184)). + Drop official support for old CUDA < 9.2 ([#887](#887)) Improved performance additions: + Reuse tmp storage in reductions in solvers and add a mutable workspace to all solvers ([#1013](#1013), [#1028](#1028)) + Add HIP unsafe atomic option for AMD ([#1091](#1091)) + Prefer vendor implementations for Dense dot, conj_dot and norm2 when available ([#967](#967)). + Tuned OpenMP SellP, COO, and ELL SpMV kernels for a small number of RHS ([#809](#809)) Fixes: + Fix various compilation warnings ([#1076](#1076), [#1183](#1183), [#1189](#1189)) + Fix issues with hwloc-related tests ([#1074](#1074)) + Fix include headers for GCC 12 ([#1071](#1071)) + Fix for simple-solver-logging example ([#1066](#1066)) + Fix for potential memory leak in Logger ([#1056](#1056)) + Fix logging of mixin classes ([#1037](#1037)) + Improve value semantics for LinOp types, like moved-from state in cross-executor copy/clones ([#753](#753)) + Fix some matrix SpMV and conversion corner cases ([#905](#905), [#978](#978)) + Fix uninitialized data ([#958](#958)) + Fix CUDA version requirement for cusparseSpSM ([#953](#953)) + Fix several issues within bash-script ([#1016](#1016)) + Fixes for `NVHPC` compiler support ([#1194](#1194)) Other additions: + Simplify and properly name GMRES kernels ([#861](#861)) + Improve pkg-config support for non-CMake libraries ([#923](#923), [#1109](#1109)) + Improve gdb pretty printer ([#987](#987), [#1114](#1114)) + Add a logger highlighting inefficient allocation and copy patterns ([#1035](#1035)) + Improved and optimized test random matrix generation ([#954](#954), [#1032](#1032)) + Better CSR strategy defaults ([#969](#969)) + Add `move_from` to `PolymorphicObject` ([#997](#997)) + Remove unnecessary device_guard usage ([#956](#956)) + Improvements to the generic accessor for mixed-precision ([#727](#727)) + Add a naive lower triangular solver implementation for CUDA ([#764](#764)) + Add support for int64 indices from CUDA 11 onward with SpMV and SpGEMM ([#897](#897)) + Add a L1 norm implementation ([#900](#900)) + Add reduce_add for arrays ([#831](#831)) + Add utility to simplify Dense View creation from an existing Dense vector ([#1136](#1136)). + Add a custom transpose implementation for Fbcsr and Csr transpose for unsupported vendor types ([#1123](#1123)) + Make IDR random initilization deterministic ([#1116](#1116)) + Move the algorithm choice for triangular solvers from Csr::strategy_type to a factory parameter ([#1088](#1088)) + Update CUDA archCoresPerSM ([#1175](#1116)) + Add kernels for Csr sparsity pattern lookup ([#994](#994)) + Differentiate between structural and numerical zeros in Ell/Sellp ([#1027](#1027)) + Add a binary IO format for matrix data ([#984](#984)) + Add a tuple zip_iterator implementation ([#966](#966)) + Simplify kernel stubs and declarations ([#888](#888)) + Simplify GKO_REGISTER_OPERATION with lambdas ([#859](#859)) + Simplify copy to device in tests and examples ([#863](#863)) + More verbose output to array assertions ([#858](#858)) + Allow parallel compilation for Jacobi kernels ([#871](#871)) + Change clang-format pointer alignment to left ([#872](#872)) + Various improvements and fixes to the benchmarking framework ([#750](#750), [#759](#759), [#870](#870), [#911](#911), [#1033](#1033), [#1137](#1137)) + Various documentation improvements ([#892](#892), [#921](#921), [#950](#950), [#977](#977), [#1021](#1021), [#1068](#1068), [#1069](#1069), [#1080](#1080), [#1081](#1081), [#1108](#1108), [#1153](#1153), [#1154](#1154)) + Various CI improvements ([#868](#868), [#874](#874), [#884](#884), [#889](#889), [#899](#899), [#903](#903), [#922](#922), [#925](#925), [#930](#930), [#936](#936), [#937](#937), [#958](#958), [#882](#882), [#1011](#1011), [#1015](#1015), [#989](#989), [#1039](#1039), [#1042](#1042), [#1067](#1067), [#1073](#1073), [#1075](#1075), [#1083](#1083), [#1084](#1084), [#1085](#1085), [#1139](#1139), [#1178](#1178), [#1187](#1187))

upsj added the 1:ST:WIP This PR is a work in progress. Not ready for review. label Jul 13, 2021

upsj self-assigned this Jul 13, 2021

ginkgo-bot added mod:cuda This is related to the CUDA module. mod:dpcpp This is related to the DPC++ module. mod:hip This is related to the HIP module. mod:openmp This is related to the OpenMP module. reg:testing This is related to testing. labels Jul 13, 2021

upsj force-pushed the simple_kernel_reduction branch from 04b20f2 to f0dbc05 Compare July 15, 2021 19:39

upsj force-pushed the simple_kernel_reduction branch from 657c5e5 to fecc953 Compare July 22, 2021 21:05

upsj added this to the Ginkgo 1.4.0 milestone Jul 24, 2021

upsj force-pushed the simple_kernel_reduction branch 2 times, most recently from 28b3ca7 to a99daf0 Compare July 27, 2021 22:36

upsj modified the milestones: Ginkgo 1.4.0, Ginkgo 1.5.0 Aug 5, 2021

upsj force-pushed the simple_kernel_reduction branch from a99daf0 to 4126f18 Compare August 10, 2021 04:47

upsj added 1:ST:ready-for-review This PR is ready for review and removed 1:ST:WIP This PR is a work in progress. Not ready for review. labels Aug 11, 2021

upsj force-pushed the simple_kernel_reduction branch 2 times, most recently from b241e2b to 4241f8a Compare August 11, 2021 10:39

upsj requested a review from a team August 11, 2021 14:16

upsj added this to Awaiting Review in Ginkgo development Aug 22, 2021

yhmtsai requested changes Aug 26, 2021

View reviewed changes

upsj mentioned this pull request Sep 14, 2021

Add device_matrix_data and device-side matrix::read #886

Merged

9 tasks

thoasm approved these changes Sep 17, 2021

View reviewed changes

upsj force-pushed the simple_kernel_reduction branch from 4241f8a to a5743fc Compare October 4, 2021 10:49

yhmtsai requested changes Oct 4, 2021

View reviewed changes

upsj and others added 20 commits October 12, 2021 18:17

add row and column OpenMP reduction kernels

98bc7df

fix types and tests

41d92b4

add dense reduction kernels

5139726

fix dpcpp simple kernel indexing

5b07de5

add reduction kernels

e9ee666

improve simple reductions for CUDA

9239139

move CUDA reduction kernels to common entirely

45202f0

add HIP kernels

879c678

fix overflows in reduction tests

6906818

avoid duplicate writes in reductions

ddbe7fd

add DPC++ reduction kernels

7103627

fix include guards for kernel launch reduction

7eec0c5

run header formatting on common kernels

53ba972

work around Intel compiler bug

5889bbc

icpc has issues with some combination of pragma omp parallel for and lambdas called inside the loop: internal error: assertion failed: find_assoc_pragma: pragma not found Putting the entire scope into an immediately evaluated lambda helps

review updates

00e6110

* remove unnecessary shmem init * add test comments

rename init to identity in reduction interface

c0093c8

simplify DPCPP kernel selection

d076d5f

simplify dpcpp local memory usage

ceaa76b

Format files

817d1d2

Co-authored-by: Tobias Ribizel <[email protected]>

fix omp name collision

dfee616

upsj force-pushed the simple_kernel_reduction branch from 70c3f3c to dfee616 Compare October 12, 2021 16:17

upsj merged commit 83bd858 into develop Oct 13, 2021

upsj deleted the simple_kernel_reduction branch October 13, 2021 08:36

upsj mentioned this pull request Oct 13, 2021

L1 norm implementation #900

Merged

greole added a commit that referenced this pull request Oct 21, 2021

Base l1 norm on #833

bc68533

greole added a commit that referenced this pull request Nov 3, 2021

Base l1 norm on #833

da7d75e

kyrillh pushed a commit to kyrillh/ginkgo that referenced this pull request Dec 17, 2021

Base l1 norm on ginkgo-project#833

36661c5

tcojean pushed a commit that referenced this pull request Nov 12, 2022

Base l1 norm on #833

b5b2c37

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Simple kernel reduction #833

Simple kernel reduction #833

upsj commented Jul 13, 2021 •

edited

Loading

sonarcloud bot commented Aug 13, 2021

codecov bot commented Aug 13, 2021 •

edited

Loading

yhmtsai left a comment

yhmtsai Aug 26, 2021

upsj Aug 26, 2021

yhmtsai Oct 4, 2021

upsj Oct 5, 2021

yhmtsai Aug 26, 2021

upsj Aug 26, 2021

thoasm left a comment

yhmtsai left a comment

yhmtsai Oct 4, 2021

upsj Oct 4, 2021

yhmtsai Oct 4, 2021

yhmtsai Oct 4, 2021

upsj Oct 5, 2021

yhmtsai Oct 4, 2021

upsj Oct 5, 2021

yhmtsai Oct 4, 2021

yhmtsai Oct 4, 2021

yhmtsai Oct 4, 2021

upsj Oct 5, 2021

yhmtsai Oct 5, 2021

yhmtsai Oct 4, 2021

upsj commented Oct 4, 2021

yhmtsai commented Oct 4, 2021

sonarcloud bot commented Oct 13, 2021

		const auto row = idx % rows;
		const auto col_block = idx / rows;

Simple kernel reduction #833

Simple kernel reduction #833

Conversation

upsj commented Jul 13, 2021 • edited Loading

sonarcloud bot commented Aug 13, 2021

codecov bot commented Aug 13, 2021 • edited Loading

Codecov Report

yhmtsai left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

thoasm left a comment

Choose a reason for hiding this comment

yhmtsai left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

upsj commented Oct 4, 2021

yhmtsai commented Oct 4, 2021

sonarcloud bot commented Oct 13, 2021

upsj commented Jul 13, 2021 •

edited

Loading

codecov bot commented Aug 13, 2021 •

edited

Loading