Minor fixes

Minor fixes

Minor fixes

example/62_conv_fwd_activ/convnd_fwd_xdl_scaleadd_scaleadd_relu_fp16.cpp

@@ -226,14 +226,17 @@ bool run_grouped_conv_fwd(bool do_verification,
 
     if(do_verification)
     {
-        auto ref_conv = ck::tensor_operation::host::ReferenceConvFwd<NDimSpatial,
-                                                                     InDataType,
-                                                                     WeiDataType,
-                                                                     OutDataType,
-                                                                     InElementOp,
-                                                                     WeiElementOp,
-                                                                     OutElementOp,
-                                                                     NumDs>();
+        auto ref_conv =
+            ck::tensor_operation::host::ReferenceConvFwd<NDimSpatial,
+                                                         InDataType,
+                                                         WeiDataType,
+                                                         OutDataType,
+                                                         InElementOp,
+                                                         WeiElementOp,
+                                                         OutElementOp,
+                                                         0, /*Num A Elementwise Tensors*/
+                                                         0, /*Num B Elementwise Tensors*/
+                                                         NumDs>();
 
         auto ref_invoker  = ref_conv.MakeInvoker();
         auto ref_argument = ref_conv.MakeArgument(in,
@@ -246,6 +249,8 @@ bool run_grouped_conv_fwd(bool do_verification,
                                                   in_element_op,
                                                   wei_element_op,
                                                   out_element_op,
+                                                  {},
+                                                  {},
                                                   d_tensors);
 
         ref_invoker.Run(ref_argument);

include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_xdl_cshuffle.hpp

@@ -104,7 +104,7 @@ __global__ void
 
     const long_index_t e_batch_offset = __builtin_amdgcn_readfirstlane(
         static_cast<long_index_t>(compute_ptr_offset_of_batch.GetEPtrOffset(g_idx)));
-    const auto ds_batch_offset = compute_ptr_offset_of_batch.GetDsPtrOffset(g_idx);
+    const auto& ds_batch_offset = compute_ptr_offset_of_batch.GetDsPtrOffset(g_idx);
 
     __shared__ char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];
 
@@ -121,13 +121,13 @@ __global__ void
         AsPointer p_as_grid_grp;
         BsPointer p_bs_grid_grp;
 
-        const auto as_batch_offset = compute_ptr_offset_of_batch.GetAsPtrOffset(g_idx);
+        const auto& as_batch_offset = compute_ptr_offset_of_batch.GetAsPtrOffset(g_idx);
 
         static constexpr index_t NumATensor = AGridDesc_AK0_M_AK1::Size();
         static_for<0, NumATensor, 1>{}(
             [&](auto i) { p_as_grid_grp(i) = p_as_grid[i] + as_batch_offset[i]; });
 
-        const auto bs_batch_offset = compute_ptr_offset_of_batch.GetBsPtrOffset(g_idx);
+        const auto& bs_batch_offset = compute_ptr_offset_of_batch.GetBsPtrOffset(g_idx);
 
         static constexpr index_t NumBTensor = BGridDesc_BK0_N_BK1::Size();
         static_for<0, NumBTensor, 1>{}(
@@ -988,7 +988,7 @@ struct DeviceGroupedConvFwdMultipleD_Xdl_CShuffle
     static auto MakeArgument(
         APointers p_as,
         BPointers p_bs,
-        std::array<const void*, NumDTensor>& p_ds,
+        const std::array<const void*, NumDTensor>& p_ds,
         void* p_e,
         const std::array<index_t, NDimSpatial + 3>& a_g_n_c_wis_lengths,
         const std::array<index_t, NDimSpatial + 3>& a_g_n_c_wis_strides,

include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_utils.hpp

@@ -20,17 +20,12 @@ struct ComputePtrOffsetOfStridedBatch
                                    index_t BatchStrideB,
                                    Array<ck::index_t, NumDTensor> BatchStrideDs,
                                    index_t BatchStrideE)
-        : BatchStrideA_(),
-          BatchStrideB_(),
+        : BatchStrideA_(BatchStrideA),
+          BatchStrideB_(BatchStrideB),
           BatchStrideDs_(BatchStrideDs),
           BatchStrideE_(BatchStrideE)
     {
-        if constexpr(!isMultiAB)
-        {
-            BatchStrideA_ = BatchStrideA;
-            BatchStrideB_ = BatchStrideB;
-        }
-        else
+        if constexpr(isMultiAB)
         {
             static_assert("Invalid constructor for multiple A or B");
         }
@@ -40,17 +35,12 @@ struct ComputePtrOffsetOfStridedBatch
                                    Array<ck::index_t, NumBTensor> BatchStrideBs,
                                    Array<ck::index_t, NumDTensor> BatchStrideDs,
                                    index_t BatchStrideE)
-        : BatchStrideA_(),
-          BatchStrideB_(),
+        : BatchStrideA_(BatchStrideAs),
+          BatchStrideB_(BatchStrideBs),
           BatchStrideDs_(BatchStrideDs),
           BatchStrideE_(BatchStrideE)
     {
-        if constexpr(isMultiAB)
-        {
-            BatchStrideA_ = BatchStrideAs;
-            BatchStrideB_ = BatchStrideBs;
-        }
-        else
+        if constexpr(!isMultiAB)
         {
             static_assert("Invalid constructor for single A and B");
         }