review updates

- make volatile load/store SFINAE more readable - add missing includes - disable assertions for HIP Co-authored-by: Thomas Grützmacher <thomas.gruetzmacher@kit.edu>
ginkgo-project · Jul 22, 2022 · 81f0a5d · 81f0a5d
1 parent 81f8241
commit 81f0a5d
Show file tree

Hide file tree

Showing 4 changed files with 10 additions and 4 deletions.
diff --git a/common/cuda_hip/components/volatile.hpp.inc b/common/cuda_hip/components/volatile.hpp.inc
@@ -50,9 +50,9 @@ load(const thrust::complex<ValueType>* values, IndexType index)
 }
 
 template <typename ValueType, typename IndexType>
-__device__ __forceinline__ void store(
-    ValueType* values, IndexType index,
-    std::enable_if_t<!is_complex_s<ValueType>::value, ValueType> value)
+__device__ __forceinline__
+    std::enable_if_t<!is_complex_s<ValueType>::value, void>
+    store(ValueType* values, IndexType index, ValueType value)
 {
     volatile ValueType* val = values + index;
     *val = value;

diff --git a/common/cuda_hip/factorization/lu_kernels.hpp.inc b/common/cuda_hip/factorization/lu_kernels.hpp.inc
@@ -119,7 +119,7 @@ __global__ __launch_bounds__(default_block_size) void factorize(
         auto val = vals[lower_nz];
         const auto diag_idx = diag_idxs[dep];
         const auto dep_end = row_ptrs[dep + 1];
-        assert(dep < row);
+        // assert(dep < row);
         if (dep_block == block) {
             // wait for a local dependency
             while (!load(sh_ready, dep_local)) {

diff --git a/reference/test/factorization/lu_kernels.cpp b/reference/test/factorization/lu_kernels.cpp
@@ -31,12 +31,15 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 ******************************<GINKGO LICENSE>*******************************/
 
 #include <algorithm>
+#include <fstream>
 #include <memory>
 
 
 #include <gtest/gtest.h>
 
 
+#include <ginkgo/core/base/array.hpp>
+#include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/csr.hpp>
 
 

diff --git a/test/factorization/lu_kernels.cpp b/test/factorization/lu_kernels.cpp
@@ -34,13 +34,16 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 
 #include <algorithm>
+#include <fstream>
 #include <memory>
 
 
 #include <gtest/gtest.h>
 
 
+#include <ginkgo/core/base/array.hpp>
 #include <ginkgo/core/base/exception.hpp>
+#include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/csr.hpp>
 #include <ginkgo/core/matrix/sparsity_csr.hpp>