address Hugh's comments

codeplaysoftware · Jan 31, 2024 · 63fcb6e · 63fcb6e
1 parent d28a70a
commit 63fcb6e
Show file tree

Hide file tree

Showing 6 changed files with 14 additions and 17 deletions.
diff --git a/src/portfft/committed_descriptor.hpp b/src/portfft/committed_descriptor.hpp
@@ -33,7 +33,7 @@
 namespace portfft {
 
 template <typename Scalar, domain Domain>
-class committed_descriptor : private committed_descriptor_impl<Scalar, Domain> {
+class committed_descriptor : private detail::committed_descriptor_impl<Scalar, Domain> {
  public:
   /**
    * Alias for `Scalar`.
@@ -46,9 +46,9 @@ class committed_descriptor : private committed_descriptor_impl<Scalar, Domain> {
   using complex_type = std::complex<Scalar>;
 
   // Use base class constructor
-  using committed_descriptor_impl<Scalar, Domain>::committed_descriptor_impl;
+  using detail::committed_descriptor_impl<Scalar, Domain>::committed_descriptor_impl;
   // Use base class function without this->
-  using committed_descriptor_impl<Scalar, Domain>::dispatch_direction;
+  using detail::committed_descriptor_impl<Scalar, Domain>::dispatch_direction;
 
   /**
    * Computes in-place forward FFT, working on a buffer.
@@ -312,4 +312,4 @@ class committed_descriptor : private committed_descriptor_impl<Scalar, Domain> {
 
 }  // namespace portfft
 
-#endif
+#endif
diff --git a/src/portfft/committed_descriptor_impl.hpp b/src/portfft/committed_descriptor_impl.hpp
@@ -39,10 +39,13 @@
 namespace portfft {
 
 template <typename Scalar, domain Domain>
-class committed_descriptor_impl;
+struct descriptor;
 
 namespace detail {
 
+template <typename Scalar, domain Domain>
+class committed_descriptor_impl;
+
 template <typename Scalar, domain Domain, detail::layout LayoutIn, detail::layout LayoutOut, Idx SubgroupSize,
           typename TIn>
 std::vector<sycl::event> compute_level(
@@ -135,12 +138,6 @@ detail::layout get_layout(const Descriptor& desc, direction dir) {
   return detail::layout::UNPACKED;
 }
 
-}  // namespace detail
-
-// forward declaration
-template <typename Scalar, domain Domain>
-struct descriptor;
-
 /*
 Compute functions in the `committed_descriptor_impl` call `dispatch_kernel` and `dispatch_kernel_helper`. These two
 functions ensure the kernel is run with a supported subgroup size. Next `dispatch_kernel_helper` calls `run_kernel`. The
@@ -1296,6 +1293,7 @@ class committed_descriptor_impl {
   }
 };
 
+}  // namespace detail
 }  // namespace portfft
 
 #endif
diff --git a/src/portfft/dispatcher/global_dispatcher.hpp b/src/portfft/dispatcher/global_dispatcher.hpp
@@ -104,8 +104,6 @@ inline IdxGlobal increment_twiddle_offset(detail::level level, Idx factor_size)
   return 0;
 }
 
-}  // namespace detail
-
 template <typename Scalar, domain Domain>
 template <typename Dummy>
 struct committed_descriptor_impl<Scalar, Domain>::calculate_twiddles_struct::inner<detail::level::GLOBAL, Dummy> {
@@ -404,6 +402,7 @@ struct committed_descriptor_impl<Scalar, Domain>::run_kernel_struct<LayoutIn, La
   }
 };
 
-}  // namespace portfft
+}  // namespace detail
+}  // namespace portfft 
 
 #endif
diff --git a/src/portfft/dispatcher/subgroup_dispatcher.hpp b/src/portfft/dispatcher/subgroup_dispatcher.hpp
@@ -583,7 +583,6 @@ PORTFFT_INLINE void subgroup_impl(const T* input, T* output, const T* input_imag
   }
   global_data.log_message_global(__func__, "exited");
 }
-}  // namespace detail
 
 template <typename Scalar, domain Domain>
 template <typename Dummy>
@@ -718,6 +717,7 @@ struct committed_descriptor_impl<Scalar, Domain>::num_scalars_in_local_mem_struc
   }
 };
 
+}  // namespace detail
 }  // namespace portfft
 
 #endif  // PORTFFT_DISPATCHER_SUBGROUP_DISPATCHER_HPP
diff --git a/src/portfft/dispatcher/workgroup_dispatcher.hpp b/src/portfft/dispatcher/workgroup_dispatcher.hpp
@@ -272,7 +272,6 @@ PORTFFT_INLINE void workgroup_impl(const T* input, T* output, const T* input_ima
   }
   global_data.log_message_global(__func__, "exited");
 }
-}  // namespace detail
 
 template <typename Scalar, domain Domain>
 template <detail::layout LayoutIn, detail::layout LayoutOut, Idx SubgroupSize, typename TIn, typename TOut>
@@ -435,6 +434,7 @@ struct committed_descriptor_impl<Scalar, Domain>::calculate_twiddles_struct::inn
   }
 };
 
+}  // namespace detail
 }  // namespace portfft
 
 #endif  // PORTFFT_DISPATCHER_WORKGROUP_DISPATCHER_HPP
diff --git a/src/portfft/dispatcher/workitem_dispatcher.hpp b/src/portfft/dispatcher/workitem_dispatcher.hpp
@@ -274,7 +274,6 @@ PORTFFT_INLINE void workitem_impl(const T* input, T* output, const T* input_imag
   }
   global_data.log_message_global(__func__, "exited");
 }
-}  // namespace detail
 
 template <typename Scalar, domain Domain>
 template <detail::layout LayoutIn, detail::layout LayoutOut, Idx SubgroupSize, typename TIn, typename TOut>
@@ -369,6 +368,7 @@ struct committed_descriptor_impl<Scalar, Domain>::calculate_twiddles_struct::inn
   }
 };
 
+}  // namespace detail
 }  // namespace portfft
 
 #endif  // PORTFFT_DISPATCHER_WORKITEM_DISPATCHER_HPP