Commit 60afb522 authored by Astha Rai's avatar Astha Rai
Browse files

removed debug print outs from device op

parent c7b0b6e7
...@@ -263,7 +263,6 @@ __global__ void ...@@ -263,7 +263,6 @@ __global__ void
} // namespace } // namespace
#ifdef CK_CODE_GEN_RTC #ifdef CK_CODE_GEN_RTC
template <typename T> template <typename T>
using is_tuple = decltype(ck::declval<T&>().IsTuple()); using is_tuple = decltype(ck::declval<T&>().IsTuple());
...@@ -753,8 +752,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle ...@@ -753,8 +752,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if(!(X == 1 && ConvStride == 1 && LeftPad == 0 && RightPad == 0)) if(!(X == 1 && ConvStride == 1 && LeftPad == 0 && RightPad == 0))
{ {
printf("Check 1");
printf("X: %d Left Pad: %d Right Pad %d Conv Stride: %d", X, LeftPad, RightPad, ConvStride );
return false; return false;
} }
} }
...@@ -771,8 +768,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle ...@@ -771,8 +768,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if(!(X == 1 && LeftPad == 0 && RightPad == 0)) if(!(X == 1 && LeftPad == 0 && RightPad == 0))
{ {
printf("Check 2");
printf("X: %d Left Pad: %d Right Pad %d", X, LeftPad, RightPad );
return false; return false;
} }
} }
...@@ -790,19 +785,11 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle ...@@ -790,19 +785,11 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if(!(ABlockTransferSrcVectorDim == 2 && C % ABlockTransferSrcScalarPerVector == 0)) if(!(ABlockTransferSrcVectorDim == 2 && C % ABlockTransferSrcScalarPerVector == 0))
{ {
printf("ABlockTransferSrcVectorDim: %d ABlockTransferSrcScalarPerVector: %d, C: %d", ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, C );
printf("Check 3");
return false; return false;
} }
} }
else else
{ {
printf("Check 4");
printf(" ALayout: %d %d %d %d %d %d %d %d %d",is_same_v<ALayout, ctc::G_NW_C> , is_same_v<ALayout, ctc::G_NHW_C> ,
is_same_v<ALayout, ctc::G_NDHW_C> , is_same_v<ALayout, ctc::GNWC> ,
is_same_v<ALayout, ctc::GNHWC> , is_same_v<ALayout, ctc::GNDHWC> ,
is_same_v<ALayout, ctc::NWGC> , is_same_v<ALayout, ctc::NHWGC> ,
is_same_v<ALayout, ctc::NDHWGC>);
return false; return false;
} }
...@@ -819,19 +806,11 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle ...@@ -819,19 +806,11 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if(!(BBlockTransferSrcVectorDim == 2 && C % BBlockTransferSrcScalarPerVector == 0)) if(!(BBlockTransferSrcVectorDim == 2 && C % BBlockTransferSrcScalarPerVector == 0))
{ {
printf("Check 5");
printf("BBlockTransferSrcVectorDim: %d BBlockTransferSrcScalarPerVector: %d, C: %d", BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, C );
return false; return false;
} }
} }
else else
{ {
printf("Check 6");
printf("B Layout: %d %d %d %d %d %d %d %d %d", is_same_v<BLayout, ctc::G_K_X_C>, is_same_v<BLayout, ctc::G_K_YX_C>,
is_same_v<BLayout, ctc::G_K_ZYX_C>, is_same_v<BLayout, ctc::GKXC>,
is_same_v<BLayout, ctc::GKYXC>, is_same_v<BLayout, ctc::GKZYXC>,
is_same_v<BLayout, ctc::KXGC>, is_same_v<BLayout, ctc::KYXGC>,
is_same_v<BLayout, ctc::KZYXGC>);
return false; return false;
} }
...@@ -851,8 +830,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle ...@@ -851,8 +830,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if(!(K % CDEBlockTransferScalarPerVector_NPerBlock == 0)) if(!(K % CDEBlockTransferScalarPerVector_NPerBlock == 0))
{ {
printf("Check 7");
printf("CDEBlockTransferScalarPerVector_NPerBlock: %d, K: %d", CDEBlockTransferScalarPerVector_NPerBlock, K );
valid = false; valid = false;
} }
...@@ -862,9 +839,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle ...@@ -862,9 +839,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if(arg.ds_g_n_k_wos_lengths_[i][0] != arg.e_g_n_k_wos_lengths_[0] || if(arg.ds_g_n_k_wos_lengths_[i][0] != arg.e_g_n_k_wos_lengths_[0] ||
arg.ds_g_n_k_wos_lengths_[i][2] != arg.e_g_n_k_wos_lengths_[2]) arg.ds_g_n_k_wos_lengths_[i][2] != arg.e_g_n_k_wos_lengths_[2])
{ {
printf("Check 8");
printf("1: arg.ds_g_n_k_wos_lengths_[i][0]: %d, arg.e_g_n_k_wos_lengths_[0]: %d", arg.ds_g_n_k_wos_lengths_[i][0], arg.e_g_n_k_wos_lengths_[0]);
printf("2: arg.ds_g_n_k_wos_lengths_[i][2]: %d, arg.e_g_n_k_wos_lengths_[2]: %d", arg.ds_g_n_k_wos_lengths_[i][2], arg.e_g_n_k_wos_lengths_[2]);
valid = false; valid = false;
} }
} }
...@@ -875,8 +849,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle ...@@ -875,8 +849,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
{ {
if(arg.ds_g_n_k_wos_lengths_[i][d] != arg.e_g_n_k_wos_lengths_[d]) if(arg.ds_g_n_k_wos_lengths_[i][d] != arg.e_g_n_k_wos_lengths_[d])
{ {
printf("Check 9");
printf("1: arg.ds_g_n_k_wos_lengths_[i][d]: %d, arg.e_g_n_k_wos_lengths_[d]: %d", arg.ds_g_n_k_wos_lengths_[i][d], arg.e_g_n_k_wos_lengths_[d]);
valid = false; valid = false;
} }
} }
...@@ -884,20 +856,12 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle ...@@ -884,20 +856,12 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
} }
else else
{ {
printf("Check 10");
printf("DLayout %d %d %d %d %d %d %d %d %d %d",is_same_v<DLayout, ctc::G_NW_K>, is_same_v<DLayout, ctc::G_NHW_K>,
is_same_v<DLayout, ctc::G_NDHW_K>, is_same_v<DLayout, ctc::GNWK>,
is_same_v<DLayout, ctc::GNHWK>, is_same_v<DLayout, ctc::GNDHWK>,
is_same_v<DLayout, ctc::NWGK>, is_same_v<DLayout, ctc::NHWGK>,
is_same_v<DLayout, ctc::NDHWGK>, is_same_v<DLayout, ctc::G_K>);
valid = false; valid = false;
} }
}); });
if(!valid) if(!valid)
{ {
printf("Check 11");
printf("Ds vector access issue");
return false; return false;
} }
...@@ -912,19 +876,11 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle ...@@ -912,19 +876,11 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
if(!(K % CDEBlockTransferScalarPerVector_NPerBlock == 0)) if(!(K % CDEBlockTransferScalarPerVector_NPerBlock == 0))
{ {
printf("Check 12");
printf("CDEBlockTransferScalarPerVector_NPerBlock: %d K: %d", CDEBlockTransferScalarPerVector_NPerBlock, K);
return false; return false;
} }
} }
else else
{ {
printf("Check 13");
printf("ELayout %d %d %d %d %d %d %d %d %d", is_same_v<ELayout, ctc::G_NW_K>, is_same_v<ELayout, ctc::G_NHW_K>,
is_same_v<ELayout, ctc::G_NDHW_K>, is_same_v<ELayout, ctc::GNWK>,
is_same_v<ELayout, ctc::GNHWK>, is_same_v<ELayout, ctc::GNDHWK>,
is_same_v<ELayout, ctc::NWGK>, is_same_v<ELayout, ctc::NHWGK>,
is_same_v<ELayout, ctc::NDHWGK>);
return false; return false;
} }
...@@ -936,7 +892,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle ...@@ -936,7 +892,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
generate_tuple([&](auto) { return arg.a_grid_desc_m_k_; }, Number<NumATensor>{}); generate_tuple([&](auto) { return arg.a_grid_desc_m_k_; }, Number<NumATensor>{});
const auto bs_grid_desc_bk0_n_bk1 = const auto bs_grid_desc_bk0_n_bk1 =
generate_tuple([&](auto) { return arg.b_grid_desc_n_k_; }, Number<NumBTensor>{}); generate_tuple([&](auto) { return arg.b_grid_desc_n_k_; }, Number<NumBTensor>{});
printf("Check 14");
return GridwiseGemm::CheckValidity(as_grid_desc_ak0_m_ak1, return GridwiseGemm::CheckValidity(as_grid_desc_ak0_m_ak1,
bs_grid_desc_bk0_n_bk1, bs_grid_desc_bk0_n_bk1,
arg.ds_grid_desc_m_n_, arg.ds_grid_desc_m_n_,
...@@ -945,7 +900,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle ...@@ -945,7 +900,6 @@ struct CodegenDeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
} }
else else
{ {
printf("Check 15");
return GridwiseGemm::CheckValidity(arg.a_grid_desc_m_k_, return GridwiseGemm::CheckValidity(arg.a_grid_desc_m_k_,
arg.b_grid_desc_n_k_, arg.b_grid_desc_n_k_,
arg.ds_grid_desc_m_n_, arg.ds_grid_desc_m_n_,
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment