fix device issue, use new GKO_FACTORY_PARAMETER_*

ginkgo-project · Aug 18, 2021 · 4a46c9f · 4a46c9f
1 parent 9f3ef86
commit 4a46c9f
Show file tree

Hide file tree

Showing 4 changed files with 5 additions and 6 deletions.
diff --git a/common/cuda_hip/multigrid/amgx_pgm_kernels.hpp.inc b/common/cuda_hip/multigrid/amgx_pgm_kernels.hpp.inc
@@ -47,7 +47,7 @@ __global__ __launch_bounds__(default_block_size) void match_edge_kernel(
     }
     auto neighbor = strongest_neighbor_vals[tidx];
     if (neighbor != -1 && strongest_neighbor_vals[neighbor] == tidx &&
-        tidx < neighbor) {
+        tidx <= neighbor) {
         // Use the smaller index as agg point
         agg_vals[tidx] = tidx;
         agg_vals[neighbor] = tidx;

diff --git a/hip/preconditioner/jacobi_kernels.hip.cpp b/hip/preconditioner/jacobi_kernels.hip.cpp
@@ -72,7 +72,6 @@ constexpr int default_num_warps = 32;
 // current GPUs have at most 84 SMs)
 constexpr int default_grid_size = 32 * 32 * 128;
 
-constexpr int default_block_size = 512;
 
 #include "common/cuda_hip/preconditioner/jacobi_kernels.hpp.inc"
 

diff --git a/include/ginkgo/core/multigrid/amgx_pgm.hpp b/include/ginkgo/core/multigrid/amgx_pgm.hpp
@@ -121,15 +121,15 @@ class AmgxPgm : public EnableLinOp<AmgxPgm<ValueType, IndexType>>,
          * NVIDIA AMGX Reference Manual (October 2017, API Version 2,
          * https://github.com/NVIDIA/AMGX/blob/main/doc/AMGX_Reference.pdf).
          */
-        unsigned GKO_FACTORY_PARAMETER(max_iterations, 15u);
+        unsigned GKO_FACTORY_PARAMETER_SCALAR(max_iterations, 15u);
 
         /**
          * The maximum ratio of unassigned number, which is valid in the
          * interval 0.0 ~ 1.0. We use the same default value as NVIDIA AMGX
          * Reference Manual (October 2017, API Version 2,
          * https://github.com/NVIDIA/AMGX/blob/main/doc/AMGX_Reference.pdf).
          */
-        double GKO_FACTORY_PARAMETER(max_unassigned_ratio, 0.05);
+        double GKO_FACTORY_PARAMETER_SCALAR(max_unassigned_ratio, 0.05);
 
         /**
          * Use the deterministic assign_to_exist_agg method or not.
@@ -138,7 +138,7 @@ class AmgxPgm : public EnableLinOp<AmgxPgm<ValueType, IndexType>>,
          * from the same matrix. Otherwise, the aggregated group might be
          * different depending on the execution ordering.
          */
-        bool GKO_FACTORY_PARAMETER(deterministic, false);
+        bool GKO_FACTORY_PARAMETER_SCALAR(deterministic, false);
     };
     GKO_ENABLE_LIN_OP_FACTORY(AmgxPgm, parameters, Factory);
     GKO_ENABLE_BUILD_METHOD(Factory);

diff --git a/omp/multigrid/amgx_pgm_kernels.cpp b/omp/multigrid/amgx_pgm_kernels.cpp
@@ -78,7 +78,7 @@ void match_edge(std::shared_ptr<const OmpExecutor> exec,
         if (agg_vals[i] == -1) {
             auto neighbor = strongest_neighbor_vals[i];
             if (neighbor != -1 && strongest_neighbor_vals[neighbor] == i &&
-                i < neighbor) {
+                i <= neighbor) {
                 // Use the smaller index as agg point
                 agg_vals[i] = i;
                 agg_vals[neighbor] = i;