ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-14 19:18:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

tmp save after little cleanup

Browse Source
This commit is contained in:
Jakub Piasecki
2025-05-28 09:27:27 +00:00
parent eeb7955556
commit 5a274f95b2
3 changed files with 17 additions and 169 deletions
Download Patch File Download Diff File Expand all files Collapse all files

60
include/ck/tensor_operation/gpu/grid/gridwise_gemm_conv_preshuffle_multiple_d_xdl_cshuffle.hpp
View File

@@ -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<ThisThreadBlock,
AElementwiseOperation,
ck::tensor_operation::element_wise::PassThrough,
InMemoryDataOperationEnum::Set,
// 1, 8, 4, 64
Sequence<NSlice, PaddedCSlice, HSlice, WSlice>, // Should be Sequence<NSlice, PaddedCSlice, HSlice, WSlice>
ACHWSliceTransferThreadClusterLengths_N_C_H_W,//ABlockTransferThreadClusterLengths_AK0_M_AK1, // should be <1, 8, 4, 2>
Sequence<NSlice, PaddedCSlice, HSlice, WSlice>,
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<AK0PerBlock, MPerBlock, AK1>,
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<ThisThreadBlock,
// AElementwiseOperation,
// ck::tensor_operation::element_wise::PassThrough,
// InMemoryDataOperationEnum::Set,
// Sequence<AK0PerBlock, MPerBlock, AK1>,
// 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<ThisThreadBlock,
@@ -917,11 +870,6 @@ struct GridwiseGemmConvFwdPreshuffleMultipleD_xdl_cshuffle
TrueKPerBlock) + ((a_grid_desc_ak0_m_ak1.GetLength(I0) * a_grid_desc_ak0_m_ak1.GetLength(I2)) %
TrueKPerBlock != 0); // TrueKPerBlock ?
// if(threadIdx.x == 0) {
// printf("hasmainkblockloop: %d\n", HasMainKBlockLoop);
// printf("Running %d k block iters", num_k_block_main_loop);
// }
gridwise_gemm_pipeline.template Run<HasMainKBlockLoop>(a_grid_desc_nchw, // global
a_block1_desc_cslice_hslice_wslice, // lds1 write
a_block1_desc_nhowo_cyx, // lds1 read

119
include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v1_nchw.hpp
View File

@@ -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<float>(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<float>(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<float>(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<float>(a_block1_buf[c * 64 * 4 + (m/64 + y) * 64 + m%64 + x - 1]);
// // auto lds2 = static_cast<float>(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<float>(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<float>(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<float>(a_block1_buf[c * 64 * 4 + (m/64 + y) * 64 + m%64 + x - 1]);
// // auto lds2 = static_cast<float>(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));

7
library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/device_grouped_conv_fwd_preshuffle_xdl_instance.hpp
View File

@@ -82,12 +82,7 @@ using device_grouped_conv_fwd_preshuffle_xdl_f16_generic_instances = std::tuple<
// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsLayout, F16, PassThrough, PassThrough, PassThrough, ConvSpec, GemmMNKPadding, 1, 128, 64, 128, 32, 8, 8, 32, 32, 2, 2, 3, 64, 3, 3, true, true, S<1, 4, 1, 8>, 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<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsLayout, F16, PassThrough, PassThrough, PassThrough, ConvSpec, GemmMNKPadding, 1, 256, 128, 64, 32, 8, 8, 32, 32, 2, 1, 4, 64, 3, 3, true, true, S<1, 4, 4, 8>, 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<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsLayout, F16, PassThrough, PassThrough, PassThrough, ConvSpec, GemmMNKPadding, 1, 256, 128, 64, 32, 8, 8, 32, 32, 2, 1, 4, 64, 3, 3, true, true, S<1, 4, 4, 8>, 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<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsLayout, F16, PassThrough, PassThrough, PassThrough, ConvSpec, GemmMNKPadding, 1, 256, 128, 64, 32, 8, 8, 32, 32, 2, 1, 4, 64, 3, 3, true, true, S<1, 4, 4, 8>, 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<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsLayout, F16, PassThrough, PassThrough, PassThrough, ConvSpec, GemmMNKPadding, 1, 256, 128, 64, 32, 8, 8, 32, 32, 2, 1, 4, 64, 3, 3, true, true, S<1, 4, 4, 8>, 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<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsLayout, F16, PassThrough, PassThrough, PassThrough, ConvSpec, GemmMNKPadding, 2, 256, 128, 64, 32, 8, 8, 32, 32, 2, 1, 4, 64, 3, 3, true, true, S<1, 4, 4, 8>, 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<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsLayout, F16, PassThrough, PassThrough, PassThrough, ConvSpec, GemmMNKPadding, 1, 256, 128, 64, 32, 8, 8, 32, 32, 2, 1, 4, 64, 3, 3, true, true, S<1, 4, 4, 8>, 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<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsLayout, F16, PassThrough, PassThrough, PassThrough, ConvSpec, GemmMNKPadding, 1, 256, 128, 64, 32, 8, 8, 32, 32, 2, 1, 4, 64, 3, 3, true, true, S<1, 4, 4, 8>, 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<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsLayout, F16, PassThrough, PassThrough, PassThrough, ConvSpec, GemmMNKPadding, 1, 256, 128, 64, 64, 8, 8, 32, 32, 2, 1, 4, 64, 3, 3, true, true, S<1, 8, 4, 8>, 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