[cuda,hip] update namespaces and includes

ginkgo-project · Aug 19, 2024 · 6b6e8e2 · 6b6e8e2
1 parent c960038
commit 6b6e8e2
Show file tree

Hide file tree

Showing 9 changed files with 63 additions and 328 deletions.
diff --git a/common/cuda_hip/base/batch_multi_vector_kernels.cpp b/common/cuda_hip/base/batch_multi_vector_kernels.cpp
@@ -37,19 +37,19 @@ void scale(std::shared_ptr<const DefaultExecutor> exec,
     const auto alpha_ub = get_batch_struct(alpha);
     const auto x_ub = get_batch_struct(x);
     if (alpha->get_common_size()[1] == 1) {
-        batch_single_kernels::scale_kernel<<<num_blocks, default_block_size, 0,
-                                             exec->get_stream()>>>(
+        GKO_DEVICE_NAMESPACE::batch_single_kernels::scale_kernel<<<
+            num_blocks, default_block_size, 0, exec->get_stream()>>>(
             alpha_ub, x_ub,
             [] __device__(int row, int col, int stride) { return 0; });
     } else if (alpha->get_common_size() == x->get_common_size()) {
-        batch_single_kernels::scale_kernel<<<num_blocks, default_block_size, 0,
-                                             exec->get_stream()>>>(
+        GKO_DEVICE_NAMESPACE::batch_single_kernels::scale_kernel<<<
+            num_blocks, default_block_size, 0, exec->get_stream()>>>(
             alpha_ub, x_ub, [] __device__(int row, int col, int stride) {
                 return row * stride + col;
             });
     } else {
-        batch_single_kernels::scale_kernel<<<num_blocks, default_block_size, 0,
-                                             exec->get_stream()>>>(
+        GKO_DEVICE_NAMESPACE::batch_single_kernels::scale_kernel<<<
+            num_blocks, default_block_size, 0, exec->get_stream()>>>(
             alpha_ub, x_ub,
             [] __device__(int row, int col, int stride) { return col; });
     }
@@ -71,11 +71,11 @@ void add_scaled(std::shared_ptr<const DefaultExecutor> exec,
     const auto x_ub = get_batch_struct(x);
     const auto y_ub = get_batch_struct(y);
     if (alpha->get_common_size()[1] == 1) {
-        batch_single_kernels::add_scaled_kernel<<<
+        GKO_DEVICE_NAMESPACE::batch_single_kernels::add_scaled_kernel<<<
             num_blocks, default_block_size, 0, exec->get_stream()>>>(
             alpha_ub, x_ub, y_ub, [] __device__(int col) { return 0; });
     } else {
-        batch_single_kernels::add_scaled_kernel<<<
+        GKO_DEVICE_NAMESPACE::batch_single_kernels::add_scaled_kernel<<<
             num_blocks, default_block_size, 0, exec->get_stream()>>>(
             alpha_ub, x_ub, y_ub, [] __device__(int col) { return col; });
     }
@@ -96,9 +96,10 @@ void compute_dot(std::shared_ptr<const DefaultExecutor> exec,
     const auto x_ub = get_batch_struct(x);
     const auto y_ub = get_batch_struct(y);
     const auto res_ub = get_batch_struct(result);
-    batch_single_kernels::compute_gen_dot_product_kernel<<<
-        num_blocks, default_block_size, 0, exec->get_stream()>>>(
-        x_ub, y_ub, res_ub, [] __device__(auto val) { return val; });
+    GKO_DEVICE_NAMESPACE::batch_single_kernels::
+        compute_gen_dot_product_kernel<<<num_blocks, default_block_size, 0,
+                                         exec->get_stream()>>>(
+            x_ub, y_ub, res_ub, [] __device__(auto val) { return val; });
 }
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(
@@ -116,9 +117,10 @@ void compute_conj_dot(std::shared_ptr<const DefaultExecutor> exec,
     const auto x_ub = get_batch_struct(x);
     const auto y_ub = get_batch_struct(y);
     const auto res_ub = get_batch_struct(result);
-    batch_single_kernels::compute_gen_dot_product_kernel<<<
-        num_blocks, default_block_size, 0, exec->get_stream()>>>(
-        x_ub, y_ub, res_ub, [] __device__(auto val) { return conj(val); });
+    GKO_DEVICE_NAMESPACE::batch_single_kernels::
+        compute_gen_dot_product_kernel<<<num_blocks, default_block_size, 0,
+                                         exec->get_stream()>>>(
+            x_ub, y_ub, res_ub, [] __device__(auto val) { return conj(val); });
 }
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(
@@ -134,9 +136,8 @@ void compute_norm2(std::shared_ptr<const DefaultExecutor> exec,
     const auto num_rhs = x->get_common_size()[1];
     const auto x_ub = get_batch_struct(x);
     const auto res_ub = get_batch_struct(result);
-    batch_single_kernels::compute_norm2_kernel<<<num_blocks, default_block_size,
-                                                 0, exec->get_stream()>>>(
-        x_ub, res_ub);
+    GKO_DEVICE_NAMESPACE::batch_single_kernels::compute_norm2_kernel<<<
+        num_blocks, default_block_size, 0, exec->get_stream()>>>(x_ub, res_ub);
 }
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(
@@ -151,7 +152,7 @@ void copy(std::shared_ptr<const DefaultExecutor> exec,
     const auto num_blocks = x->get_num_batch_items();
     const auto result_ub = get_batch_struct(result);
     const auto x_ub = get_batch_struct(x);
-    batch_single_kernels::
+    GKO_DEVICE_NAMESPACE::batch_single_kernels::
         copy_kernel<<<num_blocks, default_block_size, 0, exec->get_stream()>>>(
             x_ub, result_ub);
 }

diff --git a/common/cuda_hip/base/batch_multi_vector_kernels.hpp b/common/cuda_hip/base/batch_multi_vector_kernels.hpp
@@ -34,7 +34,6 @@
 namespace gko {
 namespace kernels {
 namespace GKO_DEVICE_NAMESPACE {
-namespace batch_multi_vector {
 namespace batch_single_kernels {
 
 
@@ -320,7 +319,6 @@ __global__ __launch_bounds__(default_block_size) void copy_kernel(
 
 
 }  // namespace batch_single_kernels
-}  // namespace batch_multi_vector
 }  // namespace GKO_DEVICE_NAMESPACE
 }  // namespace kernels
 }  // namespace gko
diff --git a/common/cuda_hip/base/batch_multi_vector_kernels.hpp.inc b/common/cuda_hip/base/batch_multi_vector_kernels.hpp.inc