diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_conv_preshuffle_multiple_d_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_conv_preshuffle_multiple_d_xdl_cshuffle.hpp index 3f0ef69a56..ed1f54ea14 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_conv_preshuffle_multiple_d_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_conv_preshuffle_multiple_d_xdl_cshuffle.hpp @@ -705,34 +705,21 @@ struct GridwiseGemmConvFwdPreshuffleMultipleD_xdl_cshuffle auto a_block1_desc_nhowo_cyx = GetABlockDescriptor_A_NHoWoCYX(); -// // constexpr auto a_block_desc_chw_slice = GetABlockDescriptor_CSlice_HSlice_WSlice(); - -// // blockwise slice global -> vgpr -> lds(slice) - -// // blockwise broadcast lds -> vgpr -> lds(m,k) - -// // then change blockwise gemm to use TrueKPerBlock - - // a blockwise copy - // global -> lds1 - - // copy bigger cslice, to a multiply of 8 or sth, so all threads will have same number of elems to copy auto a_global_into_lds1_blockwise_copy = ThreadGroupTensorSliceTransfer_v4r1, // Should be Sequence - ACHWSliceTransferThreadClusterLengths_N_C_H_W,//ABlockTransferThreadClusterLengths_AK0_M_AK1, // should be <1, 8, 4, 2> + Sequence, + ACHWSliceTransferThreadClusterLengths_N_C_H_W, Sequence<0, 1, 2, 3>,//ABlockTransferThreadClusterArrangeOrder, ADataType, AComputeDataType, decltype(a_grid_desc_nchw), decltype(a_block1_desc_cslice_hslice_wslice), Sequence<0, 1, 2, 3>,//ABlockTransferSrcAccessOrder, - Sequence<0, 1, 2, 3>,//Sequence<1, 0, 2>, + Sequence<0, 1, 2, 3>, 3,//ABlockTransferSrcVectorDim, 3, ABlockTransferSrcScalarPerVector, @@ -754,7 +741,6 @@ struct GridwiseGemmConvFwdPreshuffleMultipleD_xdl_cshuffle AElementwiseOperation, ck::tensor_operation::element_wise::PassThrough, InMemoryDataOperationEnum::Set, - // 1, 8, 4, 56 Sequence, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, @@ -766,7 +752,7 @@ struct GridwiseGemmConvFwdPreshuffleMultipleD_xdl_cshuffle Sequence<1, 0, 2>, ABlockTransferSrcVectorDim, 2, - 1,//ABlockTransferSrcScalarPerVector, // probably 1 + 1,//ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1,//ABlockTransferDstScalarPerVector_AK1, 1, 1, @@ -780,39 +766,6 @@ struct GridwiseGemmConvFwdPreshuffleMultipleD_xdl_cshuffle make_multi_index(0, 0, 0), ck::tensor_operation::element_wise::PassThrough{}); - - // change into lds1 -> lds2, lds1 transformation into ak0_m_ak1 is ready now - // A matrix blockwise copy - // auto a_blockwise_copy = - // ThreadGroupTensorSliceTransfer_v4r1, - // ABlockTransferThreadClusterLengths_AK0_M_AK1, - // ABlockTransferThreadClusterArrangeOrder, - // ADataType, - // AComputeDataType, - // decltype(a_grid_desc_ak0_m_ak1), - // decltype(a_block_desc_ak0_m_ak1), - // ABlockTransferSrcAccessOrder, - // Sequence<1, 0, 2>, - // ABlockTransferSrcVectorDim, - // 2, - // ABlockTransferSrcScalarPerVector, - // ABlockTransferDstScalarPerVector_AK1, - // 1, - // 1, - // AThreadTransferSrcResetCoordinateAfterRun, - // true, - // NumGemmKPrefetchStage>( - // a_grid_desc_ak0_m_ak1, - // make_multi_index(0, m_block_data_idx_on_grid, 0), - // a_element_op, - // a_block_desc_ak0_m_ak1, - // make_multi_index(0, 0, 0), - // ck::tensor_operation::element_wise::PassThrough{}); - // B matrix blockwise copy auto b_blockwise_copy = ThreadGroupTensorSliceTransfer_v4r1(a_grid_desc_nchw, // global a_block1_desc_cslice_hslice_wslice, // lds1 write a_block1_desc_nhowo_cyx, // lds1 read diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v1_nchw.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v1_nchw.hpp index d2a19e9af7..23e32b8f3c 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v1_nchw.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v1_nchw.hpp @@ -71,70 +71,16 @@ struct GridwiseGemmPipeline_v1_nchw<1, true, true> a_blockwise_global_to_lds1_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step); b_blockwise_copy.MoveSrcSliceWindow(b_grid_desc, b_block_copy_step); - //block_sync_lds(); + a_blockwise_global_to_lds1_copy.RunWrite(a_block1_desc_nchw_slice, a_block1_buf); b_blockwise_copy.RunWrite(b_block_desc, b_block_buf); - // c_thread_buf.Clear(); - // LDS1 -> LDS2 broadcast block_sync_lds(); - a_blockwise_lds1_to_lds2_copy.RunRead(a_block1_desc_nhowo_cyx, a_block1_buf);//a_block2_buf); ?? jaki bufor + a_blockwise_lds1_to_lds2_copy.RunRead(a_block1_desc_nhowo_cyx, a_block1_buf); a_blockwise_lds1_to_lds2_copy.RunWrite(a_block2_desc_ak0_m_k1, a_block2_buf); - //block_sync_lds(); - - // if(threadIdx.x == 0) { - // printf("Next iter\n"); - // for(int c=0; c<5; ++c) { - // for(int h=0; h<2;++h) { - // for(int w=0; w<64;++w) { - // printf("Ag[%d][%d][%d]=%f\n", c, h, w, static_cast(a_grid_buf[c * 64 * 2 + h * 64 + w])); - // } - // } - // } - - // // // A[3][2][33]=0.000000 - - // for(int c=0; c<4; ++c) { - // for(int h=0; h<4;++h) { - // for(int w=0; w<64;++w) { - // printf("A[%d][%d][%d]=%f\n", c, h, w, static_cast(a_block1_buf[c * 64 * 4 + h * 64 + w])); - // } - // } - // } - - // for(int m=0; m<128; ++m) { - // for(int k0=0; k0<4; ++k0) { - // for(int k1=0;k1<8;++k1) { - // printf("A[%d][%d]=%f\n", m, k0*8+k1, static_cast(a_block2_buf[k0*8*128 + m*8 + k1])); - - // // int k = k0*8+k1; - // // int x = k%3; - // // int y = (k/3)%3; - // // int c = k/(9); - - // // A[m][k] = A[c, y+m/64,x+m%64] z oryginalnego obrazka - - // // if(x+m%64 > 0 && x+m%64 < 127) { - // // auto lds1 = static_cast(a_block1_buf[c * 64 * 4 + (m/64 + y) * 64 + m%64 + x - 1]); - // // auto lds2 = static_cast(a_block2_buf[k0*8*128 + m*8 + k1]); - // // printf("lds2[%d][%d]:%f lds1[%d][%d][%d]:%f %s\n", m, k0*8+k1, lds2, c, y + m/64, m%64 + x - 1, lds1, (lds1 > lds2)? "diff" : ""); - - // // } - - // // if(lds1 > lds2) { - // // printf("diff lds1[%d][%d][%d]:%f lds2[%d][%d]:%f\n", c, y + m/64, m%64 + x, lds1, m, k0*8+k1 - 1, lds2); - // // } - // } - // } - // } - // } - - // // Initialize C - - - + // main body if constexpr(HasMainLoop) { @@ -142,63 +88,22 @@ struct GridwiseGemmPipeline_v1_nchw<1, true, true> do { // ask bartek how to structure this pipeline - a_blockwise_global_to_lds1_copy.RunRead(a_grid_desc, a_grid_buf); // A Global -> VGPR - b_blockwise_copy.RunRead(b_grid_desc, b_grid_buf); // B Global -> VGPR + a_blockwise_global_to_lds1_copy.RunRead(a_grid_desc, a_grid_buf); // read lds1 i + 1 + b_blockwise_copy.RunRead(b_grid_desc, b_grid_buf); // read i + 1 // block_sync_lds(); - a_blockwise_global_to_lds1_copy.RunWrite(a_block1_desc_nchw_slice, a_block1_buf); // A VGPR -> LDS1 - blockwise_gemm.Run(a_block2_buf, b_block_buf, c_thread_buf); + a_blockwise_global_to_lds1_copy.RunWrite(a_block1_desc_nchw_slice, a_block1_buf); // write lds1 i + 1 + blockwise_gemm.Run(a_block2_buf, b_block_buf, c_thread_buf); // gemm i block_sync_lds(); - a_blockwise_lds1_to_lds2_copy.RunRead(a_block1_desc_nhowo_cyx, a_block1_buf); // A LDS1 -> LDS2 + a_blockwise_lds1_to_lds2_copy.RunRead(a_block1_desc_nhowo_cyx, a_block1_buf); // read lds2 i + 1 - a_blockwise_global_to_lds1_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step); - b_blockwise_copy.MoveSrcSliceWindow(b_grid_desc, b_block_copy_step); + a_blockwise_global_to_lds1_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step); // move to i+2 + b_blockwise_copy.MoveSrcSliceWindow(b_grid_desc, b_block_copy_step); // move to i+2 - b_blockwise_copy.RunWrite(b_block_desc, b_block_buf); - a_blockwise_lds1_to_lds2_copy.RunWrite(a_block2_desc_ak0_m_k1, a_block2_buf); - // block_sync_lds(); // delete later - - // if(threadIdx.x == 0) { - // printf("Next iter\n"); - - // // // A[3][2][33]=0.000000 - - // for(int c=0; c<4; ++c) { - // for(int h=0; h<4;++h) { - // for(int w=0; w<64;++w) { - // printf("A[%d][%d][%d]=%f\n", c, h, w, static_cast(a_block1_buf[c * 64 * 4 + h * 64 + w])); - // } - // } - // } - - // for(int m=0; m<128; ++m) { - // for(int k0=0; k0<4; ++k0) { - // for(int k1=0;k1<8;++k1) { - // printf("A[%d][%d]=%f\n", m, k0*8+k1, static_cast(a_block2_buf[k0*8*128 + m*8 + k1])); - - // // int k = k0*8+k1; - // // int x = k%3; - // // int y = (k/3)%3; - // // int c = k/(9); - - // // A[m][k] = A[c, y+m/64,x+m%64] z oryginalnego obrazka - - // // if(x+m%64 > 0 && x+m%64 < 127) { - // // auto lds1 = static_cast(a_block1_buf[c * 64 * 4 + (m/64 + y) * 64 + m%64 + x - 1]); - // // auto lds2 = static_cast(a_block2_buf[k0*8*128 + m*8 + k1]); - // // printf("lds2[%d][%d]:%f lds1[%d][%d][%d]:%f %s\n", m, k0*8+k1, lds2, c, y + m/64, m%64 + x - 1, lds1, (lds1 > lds2)? "diff" : ""); - - // // } - - // // if(lds1 > lds2) { - // // printf("diff lds1[%d][%d][%d]:%f lds2[%d][%d]:%f\n", c, y + m/64, m%64 + x, lds1, m, k0*8+k1 - 1, lds2); - // // } - // } - // } - // } - // } + b_blockwise_copy.RunWrite(b_block_desc, b_block_buf); // write i + 1 + a_blockwise_lds1_to_lds2_copy.RunWrite(a_block2_desc_ak0_m_k1, a_block2_buf); // write lds2 i + 1 ++i; } while(i < (num_loop - 1)); diff --git a/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/device_grouped_conv_fwd_preshuffle_xdl_instance.hpp b/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/device_grouped_conv_fwd_preshuffle_xdl_instance.hpp index c253905454..fc047ca13f 100644 --- a/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/device_grouped_conv_fwd_preshuffle_xdl_instance.hpp +++ b/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/device_grouped_conv_fwd_preshuffle_xdl_instance.hpp @@ -82,12 +82,7 @@ using device_grouped_conv_fwd_preshuffle_xdl_f16_generic_instances = std::tuple< // DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 8>, 8>, DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, // best DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, // best - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<2, 128, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 1, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, // best - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 8, 8>, S<1, 2, 0>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, // best - - - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8> // best - + DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<2, 128, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 1, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8> // best // DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 64, 1, 4>, 8>, // best // DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, // best