remove the macro dispatch eariler than rocm 6.2

yhmtsai · yhmtsai · commit a1795d85068a · 2026-04-16T00:30:36.000+02:00
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -231,7 +231,7 @@ if(GINKGO_BUILD_HIP)
         # HIP VERSION does not use the exact version number as ROCm. Need to wait for ROCm 7.1.1 to set proper range for ROCm 7.1.0
         message(
             STATUS
-            "Disable custom thrust namespace for hip 7.1 because hip does not adapt the custom namespace fully."
+            "Disable custom thrust namespace after hip 7.1 because hip does not adapt the custom namespace fully."
         )
         set(GINKGO_HIP_CUSTOM_THRUST_NAMESPACE OFF)
     else()
diff --git a/hip/base/exception.hip.cpp b/hip/base/exception.hip.cpp
@@ -1,22 +1,14 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
 #include "ginkgo/core/base/exception.hpp"
 
 #include <string>
 
-
-#if HIP_VERSION >= 50200000
 #include <hipblas/hipblas.h>
 #include <hiprand/hiprand.h>
 #include <hipsparse/hipsparse.h>
-#else
-#include <hipblas.h>
-#include <hiprand.h>
-#include <hipsparse.h>
-#endif
-
 
 #include <ginkgo/core/base/types.hpp>
 
@@ -97,10 +89,8 @@ std::string HipsparseError::get_error(int64 error_code)
     GKO_REGISTER_HIPSPARSE_ERROR(HIPSPARSE_STATUS_INTERNAL_ERROR);
     GKO_REGISTER_HIPSPARSE_ERROR(HIPSPARSE_STATUS_MATRIX_TYPE_NOT_SUPPORTED);
     GKO_REGISTER_HIPSPARSE_ERROR(HIPSPARSE_STATUS_ZERO_PIVOT);
-#if HIP_VERSION >= 50200000
     GKO_REGISTER_HIPSPARSE_ERROR(HIPSPARSE_STATUS_NOT_SUPPORTED);
     GKO_REGISTER_HIPSPARSE_ERROR(HIPSPARSE_STATUS_INSUFFICIENT_RESOURCES);
-#endif
     return "Unknown error";
 
 #undef GKO_REGISTER_HIPSPARSE_ERROR
diff --git a/hip/base/hipblas_bindings.hip.hpp b/hip/base/hipblas_bindings.hip.hpp
@@ -1,17 +1,12 @@
-// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
 #ifndef GKO_HIP_BASE_HIPBLAS_BINDINGS_HIP_HPP_
 #define GKO_HIP_BASE_HIPBLAS_BINDINGS_HIP_HPP_
 
 
-#if HIP_VERSION >= 50200000
 #include <hipblas/hipblas.h>
-#else
-#include <hipblas.h>
-#endif
-
 
 #include <ginkgo/core/base/exception_helpers.hpp>
 #include <ginkgo/core/base/executor.hpp>
diff --git a/hip/base/hipblas_handle.hpp b/hip/base/hipblas_handle.hpp
@@ -1,16 +1,12 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
 #ifndef GKO_HIP_BASE_HIPBLAS_HANDLE_HPP_
 #define GKO_HIP_BASE_HIPBLAS_HANDLE_HPP_
 
 
-#if HIP_VERSION >= 50200000
 #include <hipblas/hipblas.h>
-#else
-#include <hipblas.h>
-#endif
 
 #include <ginkgo/core/base/exception_helpers.hpp>
 
diff --git a/hip/base/hiprand_bindings.hip.hpp b/hip/base/hiprand_bindings.hip.hpp
@@ -1,17 +1,12 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
 #ifndef GKO_HIP_BASE_HIPRAND_BINDINGS_HIP_HPP_
 #define GKO_HIP_BASE_HIPRAND_BINDINGS_HIP_HPP_
 
 
-#if HIP_VERSION >= 50200000
 #include <hiprand/hiprand.h>
-#else
-#include <hiprand.h>
-#endif
-
 
 #include <ginkgo/core/base/exception_helpers.hpp>
 
diff --git a/hip/base/hipsparse_bindings.hip.hpp b/hip/base/hipsparse_bindings.hip.hpp
@@ -1,17 +1,12 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
 #ifndef GKO_HIP_BASE_HIPSPARSE_BINDINGS_HIP_HPP_
 #define GKO_HIP_BASE_HIPSPARSE_BINDINGS_HIP_HPP_
 
 
-#if HIP_VERSION >= 50200000
 #include <hipsparse/hipsparse.h>
-#else
-#include <hipsparse.h>
-#endif
-
 
 #include <ginkgo/core/base/exception_helpers.hpp>
 #include <ginkgo/core/base/executor.hpp>
diff --git a/hip/base/hipsparse_block_bindings.hip.hpp b/hip/base/hipsparse_block_bindings.hip.hpp
@@ -1,17 +1,12 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
 #ifndef GKO_HIP_BASE_HIPSPARSE_BLOCK_BINDINGS_HIP_HPP_
 #define GKO_HIP_BASE_HIPSPARSE_BLOCK_BINDINGS_HIP_HPP_
 
 
-#if HIP_VERSION >= 50200000
 #include <hipsparse/hipsparse.h>
-#else
-#include <hipsparse.h>
-#endif
-
 
 #include <ginkgo/core/base/exception_helpers.hpp>
 
diff --git a/hip/base/hipsparse_handle.hpp b/hip/base/hipsparse_handle.hpp
@@ -1,16 +1,12 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
 #ifndef GKO_HIP_BASE_HIPSPARSE_HANDLE_HPP_
 #define GKO_HIP_BASE_HIPSPARSE_HANDLE_HPP_
 
 
-#if HIP_VERSION >= 50200000
 #include <hipsparse/hipsparse.h>
-#else
-#include <hipsparse.h>
-#endif
 
 #include <ginkgo/core/base/exception_helpers.hpp>
 
diff --git a/hip/base/memory.hip.cpp b/hip/base/memory.hip.cpp
@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -65,9 +65,6 @@ void HipAllocator::deallocate(void* dev_ptr)
 }
 
 
-#if HIP_VERSION >= 50200000
-
-
 HipAsyncAllocator::HipAsyncAllocator(hipStream_t stream) : stream_{stream} {}
 
 
@@ -86,35 +83,6 @@ void HipAsyncAllocator::deallocate(void* ptr)
 }
 
 
-#else  // Fall back to regular allocation
-
-
-HipAsyncAllocator::HipAsyncAllocator(hipStream_t stream) : stream_{stream}
-{
-#if GKO_VERBOSE_LEVEL >= 1
-    std::cerr << "This version of HIP does not support hipMallocAsync, "
-                 "please use HipAllocator instead of HipAsyncAllocator.\n";
-#endif
-}
-
-
-void* HipAsyncAllocator::allocate(size_type num_bytes)
-{
-    void* ptr{};
-    GKO_ASSERT_NO_HIP_ALLOCATION_ERRORS(hipMalloc(&ptr, num_bytes), num_bytes);
-    return ptr;
-}
-
-
-void HipAsyncAllocator::deallocate(void* ptr)
-{
-    GKO_EXIT_ON_HIP_ERROR(hipFree(ptr));
-}
-
-
-#endif
-
-
 bool HipAsyncAllocator::check_environment(int device_id,
                                           hipStream_t stream) const
 {
diff --git a/hip/base/pointer_mode_guard.hip.hpp b/hip/base/pointer_mode_guard.hip.hpp
@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -8,15 +8,8 @@
 
 #include <exception>
 
-
-#if HIP_VERSION >= 50200000
 #include <hipblas/hipblas.h>
 #include <hipsparse/hipsparse.h>
-#else
-#include <hipblas.h>
-#include <hipsparse.h>
-#endif
-
 
 #include <ginkgo/core/base/exception_helpers.hpp>
 #include <ginkgo/core/base/executor.hpp>
diff --git a/hip/base/roctx.hip.cpp b/hip/base/roctx.hip.cpp
@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -8,11 +8,7 @@
 
 
 #if GINKGO_HIP_PLATFORM_HCC && GKO_HAVE_ROCTX
-#if HIP_VERSION >= 50200000
 #include <roctracer/roctx.h>
-#else
-#include <roctx.h>
-#endif
 #endif
 
 
diff --git a/hip/base/types.hip.hpp b/hip/base/types.hip.hpp
@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -10,20 +10,13 @@
 
 #include <hip/hip_complex.h>
 #include <hip/hip_fp16.h>
-
-#include <ginkgo/core/base/types.hpp>
-
-
-#if HIP_VERSION >= 50200000
 #include <hipblas/hipblas.h>
-#else
-#include <hipblas.h>
-#endif
 #include <thrust/complex.h>
 
 #include <ginkgo/core/base/bfloat16.hpp>
 #include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/matrix_data.hpp>
+#include <ginkgo/core/base/types.hpp>
 
 #include "common/cuda_hip/base/bf16_alias.hpp"
 #include "common/cuda_hip/base/runtime.hpp"
diff --git a/hip/components/memory.hip.hpp b/hip/components/memory.hip.hpp
@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -60,7 +60,6 @@ struct gcc_atomic_intrinsic_type_map<double> {
 };
 
 
-#if HIP_VERSION >= 50100000
 // These intrinsics can be found used in clang/test/SemaCUDA/atomic-ops.cu
 // in the LLVM source code
 
@@ -70,13 +69,6 @@ struct gcc_atomic_intrinsic_type_map<double> {
     __hip_atomic_store(ptr, value, memorder, scope)
 #define HIP_SCOPE_GPU __HIP_MEMORY_SCOPE_AGENT
 #define HIP_SCOPE_THREADBLOCK __HIP_MEMORY_SCOPE_WORKGROUP
-#else
-#define HIP_ATOMIC_LOAD(ptr, memorder, scope) __atomic_load_n(ptr, memorder)
-#define HIP_ATOMIC_STORE(ptr, value, memorder, scope) \
-    __atomic_store_n(ptr, value, memorder)
-#define HIP_SCOPE_GPU -1
-#define HIP_SCOPE_THREADBLOCK -1
-#endif
 
 
 /**
diff --git a/hip/matrix/fft_kernels.hip.cpp b/hip/matrix/fft_kernels.hip.cpp
@@ -6,13 +6,7 @@
 
 #include <array>
 
-
-#if HIP_VERSION >= 50200000
 #include <hipfft/hipfft.h>
-#else
-#include <hipfft.h>
-#endif
-
 
 #include <ginkgo/core/base/exception_helpers.hpp>
 #include <ginkgo/core/base/math.hpp>
diff --git a/hip/solver/common_trs_kernels.hip.hpp b/hip/solver/common_trs_kernels.hip.hpp
@@ -9,13 +9,7 @@
 #include <functional>
 #include <memory>
 
-
-#if HIP_VERSION >= 50200000
 #include <hipsparse/hipsparse.h>
-#else
-#include <hipsparse.h>
-#endif
-
 
 #include <ginkgo/core/base/exception_helpers.hpp>
 #include <ginkgo/core/base/math.hpp>
diff --git a/hip/solver/lower_trs_kernels.hip.cpp b/hip/solver/lower_trs_kernels.hip.cpp
@@ -7,13 +7,7 @@
 #include <memory>
 #include <optional>
 
-
-#if HIP_VERSION >= 50200000
 #include <hipsparse/hipsparse.h>
-#else
-#include <hipsparse.h>
-#endif
-
 
 #include <ginkgo/core/base/exception_helpers.hpp>
 #include <ginkgo/core/base/math.hpp>
diff --git a/hip/solver/upper_trs_kernels.hip.cpp b/hip/solver/upper_trs_kernels.hip.cpp
@@ -7,13 +7,7 @@
 #include <memory>
 #include <optional>
 
-
-#if HIP_VERSION >= 50200000
 #include <hipsparse/hipsparse.h>
-#else
-#include <hipsparse.h>
-#endif
-
 
 #include <ginkgo/core/base/exception_helpers.hpp>
 #include <ginkgo/core/base/math.hpp>
diff --git a/hip/test/base/exception_helpers.hip.cpp b/hip/test/base/exception_helpers.hip.cpp
@@ -1,23 +1,16 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
 #include <hip/hip_runtime.h>
-
-#include <ginkgo/core/base/exception_helpers.hpp>
-#if HIP_VERSION >= 50200000
 #include <hipblas/hipblas.h>
 #include <hiprand/hiprand.h>
 #include <hipsparse/hipsparse.h>
-#else
-#include <hipblas.h>
-#include <hiprand.h>
-#include <hipsparse.h>
-#endif
-
 
 #include <gtest/gtest.h>
 
+#include <ginkgo/core/base/exception_helpers.hpp>
+
 
 namespace {
 
diff --git a/hip/test/matrix/fft_kernels.hip.cpp b/hip/test/matrix/fft_kernels.hip.cpp

Original file line number	Diff line number	Diff line change
`@@ -231,7 +231,7 @@ if(GINKGO_BUILD_HIP)`
`231`	`231`	`# HIP VERSION does not use the exact version number as ROCm. Need to wait for ROCm 7.1.1 to set proper range for ROCm 7.1.0`
`232`	`232`	`message(`
`233`	`233`	`STATUS`
`234`		`- "Disable custom thrust namespace for hip 7.1 because hip does not adapt the custom namespace fully."`
	`234`	`+ "Disable custom thrust namespace after hip 7.1 because hip does not adapt the custom namespace fully."`
`235`	`235`	`)`
`236`	`236`	`set(GINKGO_HIP_CUSTOM_THRUST_NAMESPACE OFF)`
`237`	`237`	`else()`