revert unneeded change

codeplaysoftware · Feb 6, 2024 · f3609cf · f3609cf
1 parent dee36f5
commit f3609cf
Showing 1 changed file with 8 additions and 6 deletions.
diff --git a/src/portfft/committed_descriptor_impl.hpp b/src/portfft/committed_descriptor_impl.hpp
@@ -963,11 +963,13 @@ class committed_descriptor_impl {
     std::size_t outer_size = total_size / params.lengths.back();
     std::size_t input_stride_0 = input_strides.back();
     std::size_t output_stride_0 = output_strides.back();
+    std::size_t input_distance_0 = n_dimensions > 1 ? params.lengths.back() : input_distance;
+    std::size_t output_distance_0 = n_dimensions > 1 ? params.lengths.back() : output_distance;
 
     PORTFFT_LOG_TRACE("Dispatching the kernel for the last dimension");
     sycl::event previous_event =
         dispatch_kernel_1d(in, out, in_imag, out_imag, dependencies, params.number_of_transforms * outer_size,
-                           input_stride_0, output_stride_0, input_distance / outer_size, output_distance / outer_size,
+                           input_stride_0, output_stride_0, input_distance_0, output_distance_0,
                            input_offset, output_offset, dimensions.back(), compute_direction);
     if (n_dimensions == 1) {
       return previous_event;
@@ -1068,6 +1070,8 @@ class committed_descriptor_impl {
                                         direction compute_direction) {
     PORTFFT_LOG_FUNCTION_ENTRY();
     if (SubgroupSize == dimension_data.used_sg_size) {
+      const bool input_packed = input_distance == dimension_data.length && input_stride == 1;
+      const bool output_packed = output_distance == dimension_data.length && output_stride == 1;
       const bool input_batch_interleaved = input_distance == 1 && input_stride == n_transforms;
       const bool output_batch_interleaved = output_distance == 1 && output_stride == n_transforms;
       for (kernel_data_struct kernel_data : dimension_data.forward_kernels) {
@@ -1087,19 +1091,17 @@ class committed_descriptor_impl {
         }
       }
 
-      // UNPACKED layout is also being dispatched as PACKED layout
-      const bool is_in_place = in == out;
-      if (!input_batch_interleaved && !output_batch_interleaved) {
+      if (input_packed && output_packed) {
         return run_kernel<detail::layout::PACKED, detail::layout::PACKED, SubgroupSize>(
             in, out, in_imag, out_imag, dependencies, n_transforms, input_offset, output_offset, dimension_data,
             compute_direction);
       }
-      if (input_batch_interleaved && !output_batch_interleaved && !is_in_place) {
+      if (input_batch_interleaved && output_packed && in != out) {
         return run_kernel<detail::layout::BATCH_INTERLEAVED, detail::layout::PACKED, SubgroupSize>(
             in, out, in_imag, out_imag, dependencies, n_transforms, input_offset, output_offset, dimension_data,
             compute_direction);
       }
-      if (!input_batch_interleaved && output_batch_interleaved && !is_in_place) {
+      if (input_packed && output_batch_interleaved && in != out) {
         return run_kernel<detail::layout::PACKED, detail::layout::BATCH_INTERLEAVED, SubgroupSize>(
             in, out, in_imag, out_imag, dependencies, n_transforms, input_offset, output_offset, dimension_data,
             compute_direction);