Rename 'GridwiseCopy' as 'GridwisePermute'

include/ck/tensor_operation/gpu/device/device_permute.hpp

@@ -11,7 +11,7 @@
 #include "ck/utility/sequence.hpp"
 #include "ck/tensor_operation/gpu/device/device_base.hpp"
 #include "ck/tensor_operation/gpu/device/matrix_padder.hpp"
-#include "ck/tensor_operation/gpu/grid/gridwise_copy.hpp"
+#include "ck/tensor_operation/gpu/grid/gridwise_permute.hpp"
 #include "ck/tensor_description/tensor_descriptor_helper.hpp"
 
 #include "ck/host_utility/kernel_launch.hpp"
@@ -140,18 +140,18 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
     using InGrid1dDesc  = decltype(MakeDescriptor_N_H_W({1, 1}, {1, 1}));
     using OutGrid1dDesc = decltype(MakeDescriptor_N_H_W({1, 1}, {1, 1}));
 
-    using GridwiseCopy = GridwiseCopy<InGrid1dDesc,
-                                      OutGrid1dDesc,
-                                      InDataTypePointer,
-                                      OutDataTypePointer,
-                                      ElementwiseOperation,
-                                      BlockSize,
-                                      NPerBlock,
-                                      HPerBlock,
-                                      WPerBlock,
-                                      MPerThread,
-                                      InScalarPerVector,
-                                      OutScalarPerVector>;
+    using GridwisePermute = GridwisePermute<InGrid1dDesc,
+                                            OutGrid1dDesc,
+                                            InDataTypePointer,
+                                            OutDataTypePointer,
+                                            ElementwiseOperation,
+                                            BlockSize,
+                                            NPerBlock,
+                                            HPerBlock,
+                                            WPerBlock,
+                                            MPerThread,
+                                            InScalarPerVector,
+                                            OutScalarPerVector>;
 
     struct Argument : public BaseArgument
     {
@@ -171,7 +171,7 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
               outLengths_(outLengths),
               outStrides_(outStrides),
               elementwise_op_(elementwise_op),
-              block_2_tile_map_(GridwiseCopy::MakeDefaultBlock2TileMap(in_grid_1d_desc_))
+              block_2_tile_map_(GridwisePermute::MakeDefaultBlock2TileMap(in_grid_1d_desc_))
         {
         }
 
@@ -187,7 +187,7 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
 
         ElementwiseOperation elementwise_op_;
 
-        typename GridwiseCopy::DefaultBlock2TileMap block_2_tile_map_;
+        typename GridwisePermute::DefaultBlock2TileMap block_2_tile_map_;
     };
 
     struct Invoker : detail::InvokerBase<Invoker, Argument>
@@ -196,13 +196,13 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
         {
             const index_t grid_size = arg.block_2_tile_map_.CalculateGridSize(arg.in_grid_1d_desc_);
 
-            const auto kernel = kernel_nd_copy<GridwiseCopy,
-                                               InGrid1dDesc,
-                                               OutGrid1dDesc,
-                                               InDataTypePointer,
-                                               OutDataTypePointer,
-                                               ElementwiseOperation,
-                                               typename GridwiseCopy::DefaultBlock2TileMap>;
+            const auto kernel = kernel_nd_permute<GridwisePermute,
+                                                  InGrid1dDesc,
+                                                  OutGrid1dDesc,
+                                                  InDataTypePointer,
+                                                  OutDataTypePointer,
+                                                  ElementwiseOperation,
+                                                  typename GridwisePermute::DefaultBlock2TileMap>;
 
             float elapsed_time = launch_and_time_kernel(stream_config,
                                                         kernel,

include/ck/tensor_operation/gpu/grid/gridwise_copy.hpp → include/ck/tensor_operation/gpu/grid/gridwise_permute.hpp

@@ -89,29 +89,29 @@ struct Block2TileMap
 };
 } // namespace detail
 
-template <typename GridwiseCopy,
+template <typename GridwisePermute,
           typename InGrid1dDesc,
           typename OutGrid1dDesc,
           typename InDataTypePointer,
           typename OutDataTypePointer,
           typename ElementwiseOperation,
           typename Block2TileMap>
-__global__ void kernel_nd_copy(const InGrid1dDesc in_grid_1d_desc,
-                               const OutGrid1dDesc out_grid_1d_desc,
-                               const InDataTypePointer p_in_global,
-                               const OutDataTypePointer p_out_global,
-                               const ElementwiseOperation elementwise_op,
-                               const Block2TileMap block_2_tile_map)
+__global__ void kernel_nd_permute(const InGrid1dDesc in_grid_1d_desc,
+                                  const OutGrid1dDesc out_grid_1d_desc,
+                                  const InDataTypePointer p_in_global,
+                                  const OutDataTypePointer p_out_global,
+                                  const ElementwiseOperation elementwise_op,
+                                  const Block2TileMap block_2_tile_map)
 {
-    __shared__ char p_shared[GridwiseCopy::GetSharedMemoryNumberOfByte()];
-
-    GridwiseCopy::Run(in_grid_1d_desc,
-                      out_grid_1d_desc,
-                      p_in_global,
-                      p_out_global,
-                      p_shared,
-                      elementwise_op,
-                      block_2_tile_map);
+    __shared__ char p_shared[GridwisePermute::GetSharedMemoryNumberOfByte()];
+
+    GridwisePermute::Run(in_grid_1d_desc,
+                         out_grid_1d_desc,
+                         p_in_global,
+                         p_out_global,
+                         p_shared,
+                         elementwise_op,
+                         block_2_tile_map);
 }
 
 template <typename InGrid1dDesc,
@@ -126,7 +126,7 @@ template <typename InGrid1dDesc,
           index_t MPerThread,
           index_t InScalarPerVector,
           index_t OutScalarPerVector>
-struct GridwiseCopy
+struct GridwisePermute
 {
     static_assert(InGrid1dDesc::GetNumOfDimension() == 3 &&
                   OutGrid1dDesc::GetNumOfDimension() == 3);
@@ -300,7 +300,7 @@ struct GridwiseCopy
             decltype(in_grid_1d_desc_tranformed),
             Sequence<0, 1, 2>, // ABlockTransferSrcAccessOrder
             Sequence<0, 1, 2>, // ABlockTransferDstAccessOrder
-            2,                 // ABlockTransferSrcVectorDim
+            1,                 // ABlockTransferSrcVectorDim
             1,                 // ABlockTransferDstVectorDim
             1,                 // ABlockTransferSrcScalarPerVector
             1,                 // ABlockTransferDstScalarPerVector