diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_base.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_base.hpp index 512a019ec8..3dc5d87630 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_base.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_base.hpp @@ -30,7 +30,8 @@ template + bool TransposeC = false, + bool LdsScalarLoad = false> struct BlockwiseGemmXdlops_pipeline_base { static constexpr auto I0 = Number<0>{}; @@ -385,7 +386,7 @@ struct BlockwiseGemmXdlops_pipeline_base Sequence<1, 1, 1, KPack>, Sequence<0, 1, 2, 3>, 3, - A_K1, + LdsScalarLoad ? 1 : A_K1, A_K1>; using BThreadCopy = ThreadwiseTensorSliceTransfer_v4, Sequence<0, 1, 2, 3>, 3, - B_K1, + LdsScalarLoad ? 1 : B_K1, B_K1>; AThreadCopy a_thread_copy_; diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_selector.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_selector.hpp index 7ddb4e9b74..9f797e7081 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_selector.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_selector.hpp @@ -32,9 +32,15 @@ template + bool DirectLoad = false, + bool LdsScalarLoad = false> constexpr auto BlockGemmPipeline_Selector() { + // Supported for Direct Load and V1 + if constexpr(LdsScalarLoad) + { + static_assert(DirectLoad && BlkGemmPipelineVer == BlockGemmPipelineVersion::v1); + } if constexpr(DirectLoad) { if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v1) @@ -58,7 +64,8 @@ constexpr auto BlockGemmPipeline_Selector() NPerXDL, MRepeat, NRepeat, - KPack>{}; + KPack, + LdsScalarLoad>{}; } else if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v4) { diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v1.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v1.hpp index b6b50a3364..cfe9fe5e10 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v1.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v1.hpp @@ -187,7 +187,6 @@ struct BlockwiseGemmXdlops_pipeline_v1( a_thread_desc_.GetElementSpaceSize()); auto b_thread_buf = make_static_buffer( @@ -213,7 +212,6 @@ struct BlockwiseGemmXdlops_pipeline_v1{}([&](auto k) { static_for<0, MRepeat, 1>{}([&](auto m0) { @@ -316,14 +313,6 @@ struct BlockwiseGemmXdlops_pipeline_v1()(ik) = b_thread_buf[Number{}]; - - // if(threadIdx.x == 0) { - // printf("a: %f b: %f\n", - // static_cast(a_thread_buf[Number{}]), - // static_cast(b_thread_buf[Number{}])); - // } }); using mfma_input_type = @@ -769,7 +758,8 @@ template + index_t KPacks, + bool LdsScalarLoad = false> struct BlockwiseGemmXdlopsDirectLoad_pipeline_v1 { }; @@ -792,9 +782,9 @@ template + bool LdsScalarLoad> struct BlockwiseGemmXdlopsDirectLoad_pipeline_v1 + KPack, + LdsScalarLoad> : BlockwiseGemmXdlops_pipeline_base + KPack, + false /*TransposeC*/, + LdsScalarLoad> { using Base = BlockwiseGemmXdlops_pipeline_base; + KPack, + false /*TransposeC*/, + LdsScalarLoad>; using Base::I0; using Base::KRepeat; using Base::xdlops_gemm; diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp index 2f1b73caaf..df72bb563f 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp @@ -1092,7 +1092,7 @@ struct GridwiseGemmMultiD_xdl_cshuffle_v3 return c_shuffle_block_desc_mblock_mperblock_nblock_nperblock; } - + static constexpr bool LdsScalarLoad = DirectLoad; using BlockwiseGemmPipe = remove_cvref_t())>; + DirectLoad, + LdsScalarLoad>())>; __device__ static constexpr index_t GetSharedMemoryNumberOfByte() { @@ -1639,7 +1640,7 @@ struct GridwiseGemmMultiD_xdl_cshuffle_v3 decltype(b_block_desc_bk0_n_bk1), BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, - BBlockTransferSrcVectorDim, // enforcer earlier + BBlockTransferSrcVectorDim, // enforced earlier BBlockTransferSrcScalarPerVector>( b_grid_desc_bk0_n_bk1, make_multi_index(num_bk0_per_block * k_idx, n_block_data_idx_on_grid, 0), diff --git a/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_bwd_data/device_grouped_conv_bwd_data_xdl_v3_instance.hpp b/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_bwd_data/device_grouped_conv_bwd_data_xdl_v3_instance.hpp index dbbebcd7e0..35178310db 100644 --- a/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_bwd_data/device_grouped_conv_bwd_data_xdl_v3_instance.hpp +++ b/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_bwd_data/device_grouped_conv_bwd_data_xdl_v3_instance.hpp @@ -65,7 +65,7 @@ using device_grouped_conv_bwd_data_xdl_v3_f16_instances = std::tuple< // ##############################################| | | | | | | | | | | | | | | | | | | | | | | | | | | | Lengths_AK0_M_AK1| ArrangeOrder| | | PerVector| PerVector_AK1| | Lengths_BK0_N_BK1| ArrangeOrder| | | PerVector| PerVector_BK1| | PerShuffle| PerShuffle| _NBlock_NPerBlock| | // ##############################################| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | // generic instance - DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffleV3, S<1, 0, 2>, S<1, 0, 2>, 2, 2, 8, 1, S<4, 8, 1>, S<0, 2, 1>, S<0, 2, 1>, 1, 2, 8, 1, 1, 1, S<1, 16, 1, 4>, S<2,2,2>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v1, F16, F16, 1, 1, true> + DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffleV3, S<1, 0, 2>, S<1, 0, 2>, 2, 2, 8, 1, S<4, 2, 8>, S<0, 2, 1>, S<0, 2, 1>, 1, 2, 8, 1, 1, 1, S<1, 16, 1, 4>, S<2,2,2>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v1, F16, F16, 1, 1, true> // DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffleV3, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<8, 8, 1>, S<0, 2, 1>, S<0, 2, 1>, 1, 4, 8, 1, 1, 1, S<1, 16, 1, 4>, S<8,8,8>>, // DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffleV3, S<1, 0, 2>, S<1, 0, 2>, 2, 4, 8, 1, S<4, 16, 1>, S<0, 2, 1>, S<0, 2, 1>, 1, 4, 8, 1, 1, 1, S<1, 16, 1, 4>, S<8,8,8>>,