prepare for unification

- Add necessary switching headers
- Provide device namespace macro via compiler definitions
- Add necessary (namespace) aliases
- adapt math lib includes and namespaces
- uniformize files

accessor/cuda_hip_helper.hpp

@@ -0,0 +1,38 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_ACCESSOR_CUDA_HIP_HELPER_HPP_
+#define GKO_ACCESSOR_CUDA_HIP_HELPER_HPP_
+
+
+#include <utility>
+
+
+#ifdef GKO_COMPILING_HIP
+#include "accessor/hip_helper.hpp"
+#else  // GKO_COMPILING_CUDA
+#include "accessor/cuda_helper.hpp"
+#endif
+
+
+namespace gko {
+namespace acc {
+
+
+template <typename AccType>
+GKO_ACC_INLINE auto as_device_range(AccType&& acc)
+{
+#ifdef GKO_COMPILING_HIP
+    return as_hip_range(std::forward<AccType>(acc));
+#else  // GKO_COMPILING_CUDA
+    return as_cuda_range(std::forward<AccType>(acc));
+#endif
+}
+
+
+}  // namespace acc
+}  // namespace gko
+
+
+#endif  // GKO_ACCESSOR_CUDA_HIP_HELPER_HPP_

cmake/create_test.cmake

@@ -152,7 +152,7 @@ endfunction(ginkgo_create_cuda_test)
 ## Internal function allowing separate test name, filename and target name
 function(ginkgo_create_cuda_test_internal test_name filename test_target_name)
     add_executable(${test_target_name} ${filename})
-    target_compile_definitions(${test_target_name} PRIVATE GKO_COMPILING_CUDA)
+    target_compile_definitions(${test_target_name} PRIVATE GKO_COMPILING_CUDA GKO_DEVICE_NAMESPACE=cuda)
     if(MSVC)
         target_compile_options(${test_target_name}
             PRIVATE
@@ -180,7 +180,7 @@ endfunction(ginkgo_create_hip_test)
 function(ginkgo_create_hip_test_internal test_name filename test_target_name)
     set_source_files_properties(${filename} PROPERTIES LANGUAGE HIP)
     add_executable(${test_target_name} ${filename})
-    target_compile_definitions(${test_target_name} PRIVATE GKO_COMPILING_HIP)
+    target_compile_definitions(${test_target_name} PRIVATE GKO_COMPILING_HIP GKO_DEVICE_NAMESPACE=hip)
     ginkgo_set_test_target_properties(${test_target_name} "_hip" ${ARGN})
     ginkgo_add_test(${test_name} ${test_target_name} ${ARGN} RESOURCE_TYPE hipgpu)
 endfunction(ginkgo_create_hip_test_internal)
@@ -195,7 +195,7 @@ endfunction()
 function(ginkgo_create_omp_test_internal test_name filename test_target_name)
     ginkgo_build_test_name(${test_name} test_target_name)
     add_executable(${test_target_name} ${test_name}.cpp)
-    target_compile_definitions(${test_target_name} PRIVATE GKO_COMPILING_OMP)
+    target_compile_definitions(${test_target_name} PRIVATE GKO_COMPILING_OMP GKO_DEVICE_NAMESPACE=omp)
     target_link_libraries(${test_target_name} PRIVATE OpenMP::OpenMP_CXX)
     ginkgo_set_test_target_properties(${test_target_name} "_omp" ${ARGN})
     ginkgo_add_test(${test_name} ${test_target_name} ${ARGN} RESOURCE_TYPE cpu)
@@ -245,7 +245,7 @@ function(ginkgo_create_common_test_internal test_name exec_type exec)
         target_link_libraries(${test_target_name} PRIVATE OpenMP::OpenMP_CXX)
     endif ()
 
-    target_compile_definitions(${test_target_name} PRIVATE EXEC_TYPE=${exec_type} EXEC_NAMESPACE=${exec} GKO_COMPILING_${exec_upper})
+    target_compile_definitions(${test_target_name} PRIVATE EXEC_TYPE=${exec_type} GKO_DEVICE_NAMESPACE=${exec} GKO_COMPILING_${exec_upper})
     target_link_libraries(${test_target_name} PRIVATE ${common_test_ADDITIONAL_LIBRARIES})
     # use float for DPC++ if necessary
     if((exec STREQUAL "dpcpp") AND GINKGO_DPCPP_SINGLE_MODE)
@@ -277,13 +277,13 @@ function(ginkgo_create_common_device_test test_name)
         # need to make a separate file for this, since we can't set conflicting properties on the same file
         configure_file(${test_name}.cpp ${test_name}.cu COPYONLY)
         ginkgo_create_cuda_test_internal(${test_name}_cuda ${CMAKE_CURRENT_BINARY_DIR}/${test_name}.cu ${test_target_name}_cuda ${ARGN})
-        target_compile_definitions(${test_target_name}_cuda PRIVATE EXEC_TYPE=CudaExecutor EXEC_NAMESPACE=cuda)
+        target_compile_definitions(${test_target_name}_cuda PRIVATE EXEC_TYPE=CudaExecutor GKO_DEVICE_NAMESPACE=cuda)
     endif()
     if(GINKGO_BUILD_HIP)
         # need to make a separate file for this, since we can't set conflicting properties on the same file
         configure_file(${test_name}.cpp ${test_name}.hip.cpp COPYONLY)
         ginkgo_create_hip_test_internal(${test_name}_hip ${CMAKE_CURRENT_BINARY_DIR}/${test_name}.hip.cpp ${test_target_name}_hip ${ARGN})
-        target_compile_definitions(${test_target_name}_hip PRIVATE EXEC_TYPE=HipExecutor EXEC_NAMESPACE=hip)
+        target_compile_definitions(${test_target_name}_hip PRIVATE EXEC_TYPE=HipExecutor GKO_DEVICE_NAMESPACE=hip)
     endif()
 endfunction(ginkgo_create_common_device_test)
 

common/cuda_hip/base/blas_bindings.hpp

@@ -0,0 +1,16 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_COMMON_CUDA_HIP_BASE_BLAS_BINDINGS_HPP_
+#define GKO_COMMON_CUDA_HIP_BASE_BLAS_BINDINGS_HPP_
+
+
+#ifdef GKO_COMPILING_HIP
+#include "hip/base/hipblas_bindings.hip.hpp"
+#else  // GKO_COMPILING_CUDA
+#include "cuda/base/cublas_bindings.hpp"
+#endif
+
+
+#endif  // GKO_COMMON_CUDA_HIP_BASE_BLAS_BINDINGS_HPP_

common/cuda_hip/base/config.hpp

@@ -0,0 +1,16 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_COMMON_CUDA_HIP_BASE_CONFIG_HPP_
+#define GKO_COMMON_CUDA_HIP_BASE_CONFIG_HPP_
+
+
+#ifdef GKO_COMPILING_HIP
+#include "hip/base/config.hip.hpp"
+#else  // GKO_COMPILING_CUDA
+#include "cuda/base/config.hpp"
+#endif
+
+
+#endif  // GKO_COMMON_CUDA_HIP_BASE_CONFIG_HPP_

common/cuda_hip/base/pointer_mode_guard.hpp

@@ -0,0 +1,16 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_COMMON_CUDA_HIP_BASE_POINTER_MODE_GUARD_HPP_
+#define GKO_COMMON_CUDA_HIP_BASE_POINTER_MODE_GUARD_HPP_
+
+
+#ifdef GKO_COMPILING_HIP
+#include "hip/base/pointer_mode_guard.hip.hpp"
+#else  // GKO_COMPILING_CUDA
+#include "cuda/base/pointer_mode_guard.hpp"
+#endif
+
+
+#endif  // GKO_COMMON_CUDA_HIP_BASE_POINTER_MODE_GUARD_HPP_

common/cuda_hip/base/randlib_bindings.hpp

@@ -0,0 +1,16 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_COMMON_CUDA_HIP_BASE_RANDLIB_BINDINGS_HPP_
+#define GKO_COMMON_CUDA_HIP_BASE_RANDLIB_BINDINGS_HPP_
+
+
+#ifdef GKO_COMPILING_HIP
+#include "hip/base/hiprand_bindings.hip.hpp"
+#else  // GKO_COMPILING_CUDA
+#include "cuda/base/curand_bindings.hpp"
+#endif
+
+
+#endif  // GKO_COMMON_CUDA_HIP_BASE_RANDLIB_BINDINGS_HPP_

common/cuda_hip/base/runtime.hpp

@@ -0,0 +1,14 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_COMMON_CUDA_HIP_BASE_RUNTIME_HPP_
+#define GKO_COMMON_CUDA_HIP_BASE_RUNTIME_HPP_
+
+
+#ifdef GKO_COMPILING_HIP
+#include <hip/hip_runtime.h>
+#endif
+
+
+#endif  // GKO_COMMON_CUDA_HIP_BASE_RUNTIME_HPP_

common/cuda_hip/base/sparselib_bindings.hpp

@@ -0,0 +1,16 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_COMMON_CUDA_HIP_BASE_SPARSELIB_BINDINGS_HPP_
+#define GKO_COMMON_CUDA_HIP_BASE_SPARSELIB_BINDINGS_HPP_
+
+
+#ifdef GKO_COMPILING_HIP
+#include "hip/base/hipsparse_bindings.hip.hpp"
+#else  // GKO_COMPILING_CUDA
+#include "cuda/base/cusparse_bindings.hpp"
+#endif
+
+
+#endif  // GKO_COMMON_CUDA_HIP_BASE_SPARSELIB_BINDINGS_HPP_

common/cuda_hip/base/thrust.hpp

@@ -0,0 +1,51 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_COMMON_CUDA_HIP_BASE_THRUST_HPP_
+#define GKO_COMMON_CUDA_HIP_BASE_THRUST_HPP_
+
+
+#include <thrust/execution_policy.h>
+
+
+#include <ginkgo/config.hpp>
+#include <ginkgo/core/base/executor.hpp>
+
+
+#if defined(GKO_COMPILING_CUDA) || \
+    (defined(GKO_COMPILING_HIP) && !GINKGO_HIP_PLATFORM_HCC)
+#include <thrust/system/cuda/detail/execution_policy.h>
+#else
+#include <thrust/system/hip/detail/execution_policy.h>
+#endif
+
+
+namespace gko {
+namespace kernels {
+namespace GKO_DEVICE_NAMESPACE {
+
+
+#ifdef GKO_COMPILING_CUDA
+inline auto thrust_policy(std::shared_ptr<const CudaExecutor> exec)
+{
+    return thrust::cuda::par.on(exec->get_stream());
+}
+#else
+inline auto thrust_policy(std::shared_ptr<const HipExecutor> exec)
+{
+#if GINKGO_HIP_PLATFORM_HCC
+    return thrust::hip::par.on(exec->get_stream());
+#else
+    return thrust::cuda::par.on(exec->get_stream());
+#endif
+}
+#endif
+
+
+}  // namespace GKO_DEVICE_NAMESPACE
+}  // namespace kernels
+}  // namespace gko
+
+
+#endif  // GKO_COMMON_CUDA_HIP_BASE_THRUST_HPP_

common/cuda_hip/base/types.hpp

@@ -0,0 +1,9 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifdef GKO_COMPILING_CUDA
+#include "cuda/base/types.hpp"
+#else
+#include "hip/base/types.hip.hpp"
+#endif

common/cuda_hip/components/atomic.hpp.inc

@@ -196,3 +196,35 @@ GKO_BIND_ATOMIC_MAX(unsigned long long int);
 
 
 #undef GKO_BIND_ATOMIC_MAX
+
+
+/**
+ * @internal
+ *
+ * @note It is not 'real' complex<float> atomic add operation
+ */
+__forceinline__ __device__ thrust::complex<float> atomic_add(
+    thrust::complex<float>* __restrict__ address, thrust::complex<float> val)
+{
+    auto addr = reinterpret_cast<float*>(address);
+    // Separate to real part and imag part
+    auto real = atomic_add(addr, val.real());
+    auto imag = atomic_add(addr + 1, val.imag());
+    return {real, imag};
+}
+
+
+/**
+ * @internal
+ *
+ * @note It is not 'real' complex<double> atomic add operation
+ */
+__forceinline__ __device__ thrust::complex<double> atomic_add(
+    thrust::complex<double>* __restrict__ address, thrust::complex<double> val)
+{
+    auto addr = reinterpret_cast<double*>(address);
+    // Separate to real part and imag part
+    auto real = atomic_add(addr, val.real());
+    auto imag = atomic_add(addr + 1, val.imag());
+    return {real, imag};
+}

common/cuda_hip/components/cooperative_groups.hpp

@@ -0,0 +1,16 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_COMMON_CUDA_HIP_COMPONENTS_COOPERATIVE_GROUPS_HPP_
+#define GKO_COMMON_CUDA_HIP_COMPONENTS_COOPERATIVE_GROUPS_HPP_
+
+
+#ifdef GKO_COMPILING_HIP
+#include "hip/components/cooperative_groups.hip.hpp"
+#else  // GKO_COMPILING_CUDA
+#include "cuda/components/cooperative_groups.cuh"
+#endif
+
+
+#endif  // GKO_COMMON_CUDA_HIP_COMPONENTS_COOPERATIVE_GROUPS_HPP_

common/cuda_hip/components/format_conversion.hpp

@@ -0,0 +1,16 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_COMMON_CUDA_HIP_COMPONENTS_FORMAT_CONVERSION_HPP_
+#define GKO_COMMON_CUDA_HIP_COMPONENTS_FORMAT_CONVERSION_HPP_
+
+
+#ifdef GKO_COMPILING_HIP
+#include "hip/components/format_conversion.hip.hpp"
+#else  // GKO_COMPILING_CUDA
+#include "cuda/components/format_conversion.cuh"
+#endif
+
+
+#endif  // GKO_COMMON_CUDA_HIP_COMPONENTS_FORMAT_CONVERSION_HPP_

common/cuda_hip/components/memory.hpp

@@ -0,0 +1,16 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_COMMON_CUDA_HIP_COMPONENTS_MEMORY_HPP_
+#define GKO_COMMON_CUDA_HIP_COMPONENTS_MEMORY_HPP_
+
+
+#ifdef GKO_COMPILING_HIP
+#include "hip/components/memory.hip.hpp"
+#else  // GKO_COMPILING_CUDA
+#include "cuda/components/memory.cuh"
+#endif
+
+
+#endif  // GKO_COMMON_CUDA_HIP_COMPONENTS_MEMORY_HPP_