Fix failing matrix tests #978

upsj · 2022-02-28T14:32:40Z

I disabled a few matrix tests on purpose to keep them from failing in #904, I will address them with this PR

Fix OpenMP/Reference Fbcsr kernels
Implement CUDA missing Fbcsr kernels
Fix CUDA Fbcsr SpMV issues
Fix Csr strategy issues

omp/matrix/fbcsr_kernels.cpp

upsj · 2022-03-01T18:10:32Z

The PR is now ready to review. There is still a bug in the Dense -> Fbcsr conversion, but I'll tackle that tomorrow.

yhmtsai

LGTM in general, there are some nit/question

common/cuda_hip/matrix/csr_common.hpp.inc

common/cuda_hip/matrix/fbcsr_kernels.hpp.inc

cuda/matrix/dense_kernels.cu

yhmtsai · 2022-03-02T13:49:39Z

hip/matrix/dense_kernels.hip.cpp

+        kernel::
+            count_nonzero_blocks_per_row<<<num_blocks, default_block_size>>>(


hipLaunchKernelGGL?

See discussion on Slack: It probably makes sense to move to the chevron launch syntax in the future, so I wanted to make a first step in that direction :)

I don't recall if there's a consensus on this. It might make sense to keep hipLaunchKernelGGL. If not, perhaps it makes sense to move this implementation to common/cuda_hip?

I checked the HIP implementation, and the macros already resolve to the Chevron syntax since 2018 inside HIP: https://github.com/ROCm-Developer-Tools/hipamd/blob/250bd582540143f53dc6d61bca963b71cf5a1087/include/hip/hcc_detail/hip_runtime.h
Is that justification enough? 😄

hip/matrix/fbcsr_kernels.hip.cpp

common/cuda_hip/matrix/dense_kernels.hpp.inc

upsj · 2022-03-04T13:51:13Z

format-rebase!

ginkgo-bot · 2022-03-04T13:55:00Z

Formatting rebase introduced changes, see Artifacts here to review them

codecov · 2022-03-08T13:41:19Z

Codecov Report

Merging #978 (b52f47b) into develop (3989e24) will decrease coverage by 0.19%.
The diff coverage is 56.18%.

❗ Current head b52f47b differs from pull request most recent head 9ba1ee9. Consider uploading reports for the commit 9ba1ee9 to get more accurate results

@@             Coverage Diff             @@
##           develop     #978      +/-   ##
===========================================
- Coverage    92.37%   92.17%   -0.20%     
===========================================
  Files          495      484      -11     
  Lines        41675    40626    -1049     
===========================================
- Hits         38496    37447    -1049     
  Misses        3179     3179

Impacted Files	Coverage Δ
include/ginkgo/core/matrix/dense.hpp	`96.29% <ø> (-0.14%)`	⬇️
omp/factorization/par_ilut_kernels.cpp	`0.00% <ø> (ø)`
omp/matrix/fbcsr_kernels.cpp	`57.53% <0.00%> (-4.69%)`	⬇️
reference/factorization/par_ilut_kernels.cpp	`99.50% <ø> (ø)`
test/matrix/matrix.cpp	`0.00% <0.00%> (ø)`
core/matrix/dense.cpp	`94.52% <97.36%> (+0.63%)`	⬆️
core/base/composition.cpp	`75.36% <100.00%> (ø)`
core/factorization/par_ic.cpp	`100.00% <100.00%> (ø)`
core/factorization/par_ict.cpp	`100.00% <100.00%> (ø)`
core/factorization/par_ilut.cpp	`100.00% <100.00%> (ø)`
... and 98 more

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 3989e24...9ba1ee9. Read the comment docs.

thoasm

LGTM!

include/ginkgo/core/base/range.hpp

yhmtsai

almost on clone should not change the stride

hip/matrix/fbcsr_kernels.hip.cpp

yhmtsai · 2022-03-03T15:41:47Z

common/cuda_hip/matrix/dense_kernels.hpp.inc

+                col < num_cols && is_nonzero(source[row * stride + col]);
+        }
+        auto nonzero_mask =
+            warp.ballot(local_nonzero) | (first_block_nonzero ? 1u : 0u);


Suggested change

warp.ballot(local_nonzero) | (first_block_nonzero ? 1u : 0u);

warp.ballot(local_nonzero) | static_cast<unsigned>(first_block_nonzero);

I don't really think this improves readability. I want to make clear that I set the lowest bit based on whether the first block already had a nonzero found previously.

core/matrix/dense.cpp

yhmtsai · 2022-03-18T11:06:07Z

core/test/matrix/dense.cpp

@@ -60,8 +60,7 @@ class Dense : public ::testing::Test {
    static void assert_equal_to_original_mtx(gko::matrix::Dense<value_type>* m)
    {
        ASSERT_EQ(m->get_size(), gko::dim<2>(2, 3));
-        ASSERT_EQ(m->get_stride(), 4);
-        ASSERT_EQ(m->get_num_stored_elements(), 2 * 4);
+        ASSERT_EQ(m->get_num_stored_elements(), 2 * m->get_stride());


I think it still needs to check the stride.

I think it makes sense to not check the stride. If the stride is for performance reasons, it need not be the same as the original stride (the user might choose something non-optimal initially, say). But generally, the stride is to allow the user to access only a part of some original memory allocation. In that case, we don't need to preserve strides while copying or cloning. While moving, we still want to keep the stride, though.

yhmtsai · 2022-03-18T11:06:56Z

core/test/matrix/dense.cpp

-    this->assert_equal_to_original_mtx(
-        dynamic_cast<decltype(this->mtx.get())>(mtx_clone.get()));
+    this->assert_equal_to_original_mtx(mtx_clone.get());
+    ASSERT_EQ(mtx_clone->get_stride(), 3);


clone should not change the stride

I think this is fine - clone should ignore the original stride.

yhmtsai · 2022-03-18T13:04:16Z

common/cuda_hip/matrix/dense_kernels.hpp.inc

+        // assert(last_thread >= 0);
+        // mask off everything below first_thread
+        const auto lower_mask =
+            first_thread < 0 ? full_mask : ~((one_mask << first_thread) - 1u);


Suggested change

first_thread < 0 ? full_mask : ~((one_mask << first_thread) - 1u);

first_thread < 0 ? full_mask : full_mask - ((one_mask << first_thread) - 1u);

or use xor operation. it's more for dpcpp because dpcpp the subgroup is less than the type length

it should not really matter, since the purpose of the mask is to mask off bits from other threads. The ballot call in dpcpp should not return bits outside its own full mask.

yhmtsai · 2022-03-18T13:05:34Z

common/cuda_hip/matrix/dense_kernels.hpp.inc

+        block_count +=
+            (block_local_col == block_size - 1 && local_mask) ? 1 : 0;


Suggested change

block_count +=

(block_local_col == block_size - 1 && local_mask) ? 1 : 0;

block_count +=

static_cast<IndexType>(block_local_col == block_size - 1 && local_mask);

if want to reduce the branch

simple ternary operations usually don't get translated into branches, and this one specifically doesn't even need a cmov, because it is just a written-out conversion bool -> int for clarity.

Slaedr

LGTM! Great job completing more of the Fbcsr kernels; I see you had some fun doing that :) I think the stride handling also looks good to me.

common/cuda_hip/matrix/csr_common.hpp.inc

Slaedr · 2022-03-18T13:16:34Z

common/cuda_hip/matrix/csr_kernels.hpp.inc

+                values[i * bs * bs + (in_cols[nz] % bs) * bs +
+                       (in_rows[nz] % bs)] = fake_complex_unpack(in_vals[nz]);
+            }
+        });


Great implementation overall. If this loop could be written more over the original index space (over nnz rather than over num_blocks), we could get good performance. But I guess performance is not crucial here, so this will do.

this would require a lookup nonzero -> block index, which we don't have at that point. We can replace this by our own warp-parallel kernel in the future, where we can take care of load-balancing between blocks, and doing (almost) coalescing reads.

common/cuda_hip/matrix/dense_kernels.hpp.inc

Slaedr · 2022-03-19T22:51:03Z

core/test/matrix/dense.cpp

-    this->assert_equal_to_original_mtx(
-        dynamic_cast<decltype(this->mtx.get())>(mtx_clone.get()));
+    this->assert_equal_to_original_mtx(mtx_clone.get());
+    ASSERT_EQ(mtx_clone->get_stride(), 3);


I think this is fine - clone should ignore the original stride.

Slaedr · 2022-03-19T23:08:46Z

hip/matrix/dense_kernels.hip.cpp

+        kernel::
+            count_nonzero_blocks_per_row<<<num_blocks, default_block_size>>>(


I don't recall if there's a consensus on this. It might make sense to keep hipLaunchKernelGGL. If not, perhaps it makes sense to move this implementation to common/cuda_hip?

hip/matrix/fbcsr_kernels.hip.cpp

cuda/matrix/fbcsr_kernels.cu

core/test/matrix/dense.cpp

Slaedr

LGTM! Just once concern below.

Slaedr · 2022-04-11T15:55:53Z

core/matrix/dense.cpp

-        result->values_ = this->values_;
-        result->stride_ = this->stride_;
-        result->set_size(this->get_size());
+    if (result->get_size() != this->get_size()) {


If the original size is larger than required, perhaps we should preserve the original allocation. Then the new stride would be the original stride, while the new size is the new smaller size.

Not sure about this. It might encourage users to keep pointers to the original allocation, which might cause issues if we change this behavior in the future.

yhmtsai

LGTM (except for clone)

yhmtsai · 2022-04-19T10:44:08Z

common/cuda_hip/matrix/dense_kernels.hpp.inc

+        // only consider threads in the current block
+        const auto first_thread = block_base_col - base_col;
+        const auto last_thread = first_thread + block_size;
+        // HIP compiles these assertions in Release, traps unconditionally


This is the issue in v5.0 and fixed in v5.0.2.
related doc: https://docs.amd.com/bundle/ROCm-Release-Notes-v5.0.2/page/Fixed_Defects_in_This_Release.html

It's still broken for strided accesses though

Co-authored-by: Yuhsiang Tsai <yhmtsai@gmail.com>

This also changes the Dense convert_to(Dense) behavior so that it preserves strides only if no reallocation is necessary, or on moves

This means that padding in the target will be copied and written back after the copy has finished. This fixes issues in builds without OpenMP executor and reference/test/matrix/dense_kernels

This is only relevant for really old HIP/CUDA, but still causes test failures.

* use ConstArrayView for Dense::create_real_view * improve readability of Dense -> Fbcsr conversion * update documentation * remove unnecessary template parameters * more constexpr * fix generic matrix test formatting Co-authored-by: Aditya Kashi <aditya.kashi@kit.edu> Co-authored-by: Yuhsiang M. Tsai <yhmtsai@gmail.com>

ginkgo-bot · 2022-04-22T14:25:21Z

Note: This PR changes the Ginkgo ABI:

Functions changes summary: 4 Removed, 0 Changed, 16 Added functions
Variables changes summary: 0 Removed, 0 Changed, 0 Added variable

For details check the full ABI diff under Artifacts here

sonarqubecloud · 2022-04-23T13:10:41Z

SonarCloud Quality Gate failed.

0 Bugs
0 Vulnerabilities
0 Security Hotspots
8 Code Smells

32.8% Coverage
13.7% 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 Feb 28, 2022

upsj self-assigned this Feb 28, 2022

ginkgo-bot added mod:core This is related to the core module. mod:openmp This is related to the OpenMP module. mod:reference This is related to the reference module. reg:testing This is related to testing. type:matrix-format This is related to the Matrix formats labels Feb 28, 2022

Slaedr reviewed Feb 28, 2022

View reviewed changes

omp/matrix/fbcsr_kernels.cpp Show resolved Hide resolved

upsj force-pushed the spmv_fixes branch from b5113d4 to abf114d Compare March 1, 2022 18:09

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 Mar 1, 2022

thoasm self-requested a review March 2, 2022 10:49

yhmtsai approved these changes Mar 2, 2022

View reviewed changes

upsj force-pushed the spmv_fixes branch from 9cafe5b to 3173263 Compare March 3, 2022 11:18

ginkgo-bot force-pushed the spmv_fixes branch from 2b15e8a to 3dc4313 Compare March 4, 2022 13:55

upsj added the 1:ST:run-full-test label Mar 8, 2022

upsj requested a review from yhmtsai March 10, 2022 18:37

Slaedr self-requested a review March 17, 2022 11:14

thoasm approved these changes Mar 17, 2022

View reviewed changes

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

yhmtsai requested changes Mar 18, 2022

View reviewed changes

Slaedr approved these changes Mar 21, 2022

View reviewed changes

upsj force-pushed the spmv_fixes branch from 58cf056 to fed44bf Compare April 7, 2022 16:49

upsj requested a review from yhmtsai April 7, 2022 16:50

upsj force-pushed the spmv_fixes branch from 1991005 to 90cd235 Compare April 11, 2022 15:20

Slaedr approved these changes Apr 11, 2022

View reviewed changes

yhmtsai approved these changes Apr 19, 2022

View reviewed changes

upsj 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 Apr 20, 2022

upsj and others added 16 commits April 21, 2022 17:59

fix OpenMP Fbcsr SpMV and move-read

1e5af84

add CUDA Fbcsr -> Dense conversion

3b0030f

fix (Fb)Csr strategies and add Fbcsr conversions

1f1f276

improve and fix Dense -> Fbcsr conversion

20562f4

fix HIP Fbcsr SpMV

8ae8a94

It's still broken for strided accesses though

review updates

5dbb787

Co-authored-by: Yuhsiang Tsai <yhmtsai@gmail.com>

replace Array<...>::view by make_array_view

b57106a

add stride to convert_to(Dense) tests

f7e327d

This also changes the Dense convert_to(Dense) behavior so that it preserves strides only if no reallocation is necessary, or on moves

fix submatrix generation for empty and full rows

bcd48e5

test and fix SpMV interaction with stride/padding

5d798ea

fix DPC++ Csr merge_path SpMV segfault

592d93f

initiate Dense cross-executor copy from source

0b74a65

This means that padding in the target will be copied and written back after the copy has finished. This fixes issues in builds without OpenMP executor and reference/test/matrix/dense_kernels

fix dpcpp Csr classical SpMV

3dd7aa5

remove assertion causing hip-nvcc to emit traps

11737d7

This is only relevant for really old HIP/CUDA, but still causes test failures.

fix dpcpp Csr classical SpMV with subgroup_size 1

480bbb5

upsj force-pushed the spmv_fixes branch from 90cd235 to 1d7ac95 Compare April 21, 2022 16:03

use ROCm 4.0 instead of 3.5

9ba1ee9

upsj force-pushed the spmv_fixes branch from e32c14d to 9ba1ee9 Compare April 22, 2022 14:11

upsj merged commit 06353d3 into develop Apr 23, 2022

upsj deleted the spmv_fixes branch April 23, 2022 11:10

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Fix failing matrix tests #978

Fix failing matrix tests #978

upsj commented Feb 28, 2022 •

edited

Loading

upsj commented Mar 1, 2022

yhmtsai left a comment

yhmtsai Mar 2, 2022

upsj Mar 4, 2022

Slaedr Mar 19, 2022

upsj Apr 8, 2022

upsj commented Mar 4, 2022

ginkgo-bot commented Mar 4, 2022

codecov bot commented Mar 8, 2022 •

edited

Loading

thoasm left a comment

yhmtsai left a comment

yhmtsai Mar 3, 2022

upsj Mar 22, 2022

yhmtsai Mar 18, 2022

Slaedr Mar 19, 2022

yhmtsai Mar 18, 2022

Slaedr Mar 19, 2022 •

edited

Loading

yhmtsai Mar 18, 2022

upsj Mar 22, 2022

yhmtsai Mar 18, 2022

upsj Mar 22, 2022 •

edited

Loading

Slaedr left a comment •

edited

Loading

Slaedr Mar 18, 2022

upsj Mar 22, 2022

Slaedr Mar 19, 2022 •

edited

Loading

Slaedr Mar 19, 2022

Slaedr left a comment

Slaedr Apr 11, 2022

upsj Apr 12, 2022

yhmtsai left a comment

yhmtsai Apr 19, 2022

ginkgo-bot commented Apr 22, 2022

sonarqubecloud bot commented Apr 23, 2022

		kernel::
		count_nonzero_blocks_per_row<<<num_blocks, default_block_size>>>(

	warp.ballot(local_nonzero) \| (first_block_nonzero ? 1u : 0u);
	warp.ballot(local_nonzero) \| static_cast<unsigned>(first_block_nonzero);

	first_thread < 0 ? full_mask : ~((one_mask << first_thread) - 1u);
	first_thread < 0 ? full_mask : full_mask - ((one_mask << first_thread) - 1u);

		block_count +=
		(block_local_col == block_size - 1 && local_mask) ? 1 : 0;

Fix failing matrix tests #978

Fix failing matrix tests #978

Conversation

upsj commented Feb 28, 2022 • edited Loading

upsj commented Mar 1, 2022

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

upsj commented Mar 4, 2022

ginkgo-bot commented Mar 4, 2022

codecov bot commented Mar 8, 2022 • edited Loading

Codecov Report

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

Slaedr Mar 19, 2022 • edited Loading

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 Mar 22, 2022 • edited Loading

Choose a reason for hiding this comment

Slaedr left a comment • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Slaedr Mar 19, 2022 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Slaedr left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this 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

ginkgo-bot commented Apr 22, 2022

sonarqubecloud bot commented Apr 23, 2022

upsj commented Feb 28, 2022 •

edited

Loading

codecov bot commented Mar 8, 2022 •

edited

Loading

Slaedr Mar 19, 2022 •

edited

Loading

upsj Mar 22, 2022 •

edited

Loading

Slaedr left a comment •

edited

Loading

Slaedr Mar 19, 2022 •

edited

Loading