From a23a3dbaca747964c1a5088396fb645967e8cb7e Mon Sep 17 00:00:00 2001
From: Yuuichi Asahi <y.asahi@nr.titech.ac.jp>
Date: Wed, 8 Jan 2025 16:41:42 +0900
Subject: [PATCH 1/3] fix: buffer size for rocfft execution info

---
 fft/src/KokkosFFT_ROCM_types.hpp | 24 ++++++++++++------------
 1 file changed, 12 insertions(+), 12 deletions(-)

diff --git a/fft/src/KokkosFFT_ROCM_types.hpp b/fft/src/KokkosFFT_ROCM_types.hpp
index ed8b06a8..05856d68 100644
--- a/fft/src/KokkosFFT_ROCM_types.hpp
+++ b/fft/src/KokkosFFT_ROCM_types.hpp
@@ -66,15 +66,12 @@ struct ScopedRocfftPlanDescription {
 };
 
 /// \brief A class that wraps rocfft_execution_info for RAII
-template <typename FloatingPointType>
 struct ScopedRocfftExecutionInfo {
  private:
-  using BufferViewType =
-      Kokkos::View<Kokkos::complex<FloatingPointType> *, Kokkos::HIP>;
   rocfft_execution_info m_execution_info;
 
   //! Internal work buffer
-  BufferViewType m_buffer;
+  void *m_workbuffer = nullptr;
 
  public:
   ScopedRocfftExecutionInfo() {
@@ -84,6 +81,10 @@ struct ScopedRocfftExecutionInfo {
                        "rocfft_execution_info_create failed");
   }
   ~ScopedRocfftExecutionInfo() noexcept {
+    if (m_workbuffer != nullptr) {
+      hipError_t hip_status = hipFree(m_workbuffer);
+      if (hip_status != hipSuccess) Kokkos::abort("hipFree failed");
+    }
     rocfft_status status = rocfft_execution_info_destroy(m_execution_info);
     if (status != rocfft_status_success)
       Kokkos::abort("rocfft_execution_info_destroy failed");
@@ -111,9 +112,10 @@ struct ScopedRocfftExecutionInfo {
 
     // Set work buffer
     if (workbuffersize > 0) {
-      m_buffer = BufferViewType("workbuffer", workbuffersize);
-      status   = rocfft_execution_info_set_work_buffer(
-          m_execution_info, (void *)m_buffer.data(), workbuffersize);
+      hipError_t hip_status = hipMalloc(&m_workbuffer, workbuffersize);
+      KOKKOSFFT_THROW_IF(hip_status != hipSuccess, "hipMalloc failed");
+      status = rocfft_execution_info_set_work_buffer(
+          m_execution_info, m_workbuffer, workbuffersize);
       KOKKOSFFT_THROW_IF(status != rocfft_status_success,
                          "rocfft_execution_info_set_work_buffer failed");
     }
@@ -124,14 +126,12 @@ struct ScopedRocfftExecutionInfo {
 template <typename T>
 struct ScopedRocfftPlan {
  private:
-  using floating_point_type = KokkosFFT::Impl::base_floating_point_type<T>;
-  using ScopedRocfftExecutionInfoType =
-      ScopedRocfftExecutionInfo<floating_point_type>;
+  using floating_point_type    = KokkosFFT::Impl::base_floating_point_type<T>;
   rocfft_precision m_precision = std::is_same_v<floating_point_type, float>
                                      ? rocfft_precision_single
                                      : rocfft_precision_double;
   rocfft_plan m_plan;
-  std::unique_ptr<ScopedRocfftExecutionInfoType> m_execution_info;
+  std::unique_ptr<ScopedRocfftExecutionInfo> m_execution_info;
 
  public:
   ScopedRocfftPlan(const FFTWTransformType transform_type,
@@ -209,7 +209,7 @@ struct ScopedRocfftPlan {
     KOKKOSFFT_THROW_IF(status != rocfft_status_success,
                        "rocfft_plan_get_work_buffer_size failed");
 
-    m_execution_info = std::make_unique<ScopedRocfftExecutionInfoType>();
+    m_execution_info = std::make_unique<ScopedRocfftExecutionInfo>();
     m_execution_info->setup(exec_space, workbuffersize);
   }
 

From 109e50248dd70abdcb7e916feb8addef6f705b65 Mon Sep 17 00:00:00 2001
From: Yuuichi Asahi <y.asahi@nr.titech.ac.jp>
Date: Wed, 8 Jan 2025 18:21:51 +0900
Subject: [PATCH 2/3] use kokkos_malloc instead of hipMalloc

---
 fft/src/KokkosFFT_ROCM_types.hpp | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/fft/src/KokkosFFT_ROCM_types.hpp b/fft/src/KokkosFFT_ROCM_types.hpp
index 05856d68..ddfaea1f 100644
--- a/fft/src/KokkosFFT_ROCM_types.hpp
+++ b/fft/src/KokkosFFT_ROCM_types.hpp
@@ -82,8 +82,7 @@ struct ScopedRocfftExecutionInfo {
   }
   ~ScopedRocfftExecutionInfo() noexcept {
     if (m_workbuffer != nullptr) {
-      hipError_t hip_status = hipFree(m_workbuffer);
-      if (hip_status != hipSuccess) Kokkos::abort("hipFree failed");
+      Kokkos::kokkos_free<Kokkos::HIP>(m_workbuffer);
     }
     rocfft_status status = rocfft_execution_info_destroy(m_execution_info);
     if (status != rocfft_status_success)
@@ -112,8 +111,9 @@ struct ScopedRocfftExecutionInfo {
 
     // Set work buffer
     if (workbuffersize > 0) {
-      hipError_t hip_status = hipMalloc(&m_workbuffer, workbuffersize);
-      KOKKOSFFT_THROW_IF(hip_status != hipSuccess, "hipMalloc failed");
+      m_workbuffer = Kokkos::kokkos_malloc<Kokkos::HIP>(
+          "kokkos_malloc workbuffer", workbuffersize);
+
       status = rocfft_execution_info_set_work_buffer(
           m_execution_info, m_workbuffer, workbuffersize);
       KOKKOSFFT_THROW_IF(status != rocfft_status_success,

From 7c3084707b25e37b83bdc657b59e8bc6f820677b Mon Sep 17 00:00:00 2001
From: Yuuichi Asahi <y.asahi@nr.titech.ac.jp>
Date: Wed, 8 Jan 2025 18:28:29 +0900
Subject: [PATCH 3/3] fix annotation for kokkos_malloc

---
 fft/src/KokkosFFT_ROCM_types.hpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/fft/src/KokkosFFT_ROCM_types.hpp b/fft/src/KokkosFFT_ROCM_types.hpp
index ddfaea1f..8bf19c07 100644
--- a/fft/src/KokkosFFT_ROCM_types.hpp
+++ b/fft/src/KokkosFFT_ROCM_types.hpp
@@ -111,8 +111,8 @@ struct ScopedRocfftExecutionInfo {
 
     // Set work buffer
     if (workbuffersize > 0) {
-      m_workbuffer = Kokkos::kokkos_malloc<Kokkos::HIP>(
-          "kokkos_malloc workbuffer", workbuffersize);
+      m_workbuffer =
+          Kokkos::kokkos_malloc<Kokkos::HIP>("workbuffer", workbuffersize);
 
       status = rocfft_execution_info_set_work_buffer(
           m_execution_info, m_workbuffer, workbuffersize);