changing class to POD

composable_kernel/include/driver/driver_dynamic_convolution_forward_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp

@@ -950,6 +950,11 @@ struct DriverDynamicConvolutionForwardImplicitGemm_v4r4_nchw_kcyx_nkhw_pad
                                                           ck::Tuple<ck::Sequence<1, 2, 3, 4>>,
                                                           ck::Sequence<1, 2, 3, 4>>();
 
+        constexpr auto coord = make_dynamic_tensor_coordinate(desc, make_multi_index(0, 0, 0, 0));
+
+        constexpr auto iter = make_dynamic_tensor_coordinate_iterator(
+            desc, make_multi_index(0, 0, 0, 0), Sequence<0>{});
+
         static_assert(std::is_trivial<Sequence<1>>::value, "wrong");
         static_assert(std::is_trivial<detail::TupleElementKey<0>>::value, "wrong");
         static_assert(
@@ -970,6 +975,10 @@ struct DriverDynamicConvolutionForwardImplicitGemm_v4r4_nchw_kcyx_nkhw_pad
             std::is_trivial<
                 remove_reference_t<remove_cv_t<decltype(wei_k_c_y_x_global_desc)>>>::value,
             "wrong");
+        static_assert(std::is_trivial<remove_reference_t<remove_cv_t<decltype(coord)>>>::value,
+                      "wrong");
+        static_assert(std::is_trivial<remove_reference_t<remove_cv_t<decltype(iter)>>>::value,
+                      "wrong");
     }
 };
 #endif

composable_kernel/include/tensor_description/dynamic_tensor_descriptor.hpp

@@ -163,6 +163,8 @@ struct DynamicTensorCoordinate
     using VisibleIndex = MultiIndex<ndim_visible_>;
 
     public:
+    __host__ __device__ constexpr DynamicTensorCoordinate() = default;
+
     __host__ __device__ constexpr DynamicTensorCoordinate(const HiddenIndex& idx_hidden)
         : idx_hidden_{idx_hidden}
     {
@@ -193,6 +195,8 @@ struct DynamicTensorCoordinateIterator
     using VisibleIndex = MultiIndex<NDimVisible>;
 
     public:
+    __host__ __device__ constexpr DynamicTensorCoordinateIterator() = default;
+
     __host__ __device__ constexpr DynamicTensorCoordinateIterator(
         const VisibleIndex& idx_diff_visible, const MultiIndex<NTransform>& do_transforms)
         : idx_diff_visible_{idx_diff_visible}, do_transforms_{do_transforms}
@@ -207,8 +211,8 @@ struct DynamicTensorCoordinateIterator
         return idx_diff_visible_;
     }
 
-    const VisibleIndex idx_diff_visible_;
-    const MultiIndex<NTransform> do_transforms_;
+    VisibleIndex idx_diff_visible_;
+    MultiIndex<NTransform> do_transforms_;
 
     // HACK: control UpdateLowerIndex()
     static constexpr UpdateLowerIndexHack update_lower_index_hack_;

composable_kernel/include/tensor_operation/threadwise_dynamic_tensor_slice_transfer.hpp

@@ -327,14 +327,11 @@ struct ThreadwiseDynamicTensorSliceTransfer_v3
     static constexpr index_t nDim = SliceLengths::Size();
     using Index                   = MultiIndex<nDim>;
 
-    using SrcDesc_ = remove_reference_t<remove_cv_t<SrcDesc>>;
-    using DstDesc_ = remove_reference_t<remove_cv_t<DstDesc>>;
-
-    using SrcCoord = decltype(make_dynamic_tensor_coordinate(SrcDesc_{}, Index{}));
-    using DstCoord = decltype(make_dynamic_tensor_coordinate(DstDesc_{}, Index{}));
+    using SrcCoord = decltype(make_dynamic_tensor_coordinate(SrcDesc{}, Index{}));
+    using DstCoord = decltype(make_dynamic_tensor_coordinate(DstDesc{}, Index{}));
 
-    using SrcCoordIterator = decltype(make_dynamic_tensor_coordinate_iterator(SrcDesc_{}, Index{}));
-    using DstCoordIterator = decltype(make_dynamic_tensor_coordinate_iterator(DstDesc_{}, Index{}));
+    using SrcCoordIterator = decltype(make_dynamic_tensor_coordinate_iterator(SrcDesc{}, Index{}));
+    using DstCoordIterator = decltype(make_dynamic_tensor_coordinate_iterator(DstDesc{}, Index{}));
 
     __device__ constexpr ThreadwiseDynamicTensorSliceTransfer_v3(const SrcDesc& src_desc,
                                                                  const Index& src_slice_origin,