diff --git a/example/20_grouped_conv_bwd_weight/CMakeLists.txt b/example/20_grouped_conv_bwd_weight/CMakeLists.txt index 6fbaee7dba..de54c4a4b9 100644 --- a/example/20_grouped_conv_bwd_weight/CMakeLists.txt +++ b/example/20_grouped_conv_bwd_weight/CMakeLists.txt @@ -1,6 +1,7 @@ add_custom_target(example_grouped_conv_bwd_weight) add_example_executable(example_grouped_conv_bwd_weight_xdl_fp16 grouped_conv_bwd_weight_xdl_fp16.cpp) add_example_dependencies(example_grouped_conv_bwd_weight example_grouped_conv_bwd_weight_xdl_fp16) +target_compile_options(example_grouped_conv_bwd_weight_xdl_fp16 PRIVATE -save-temps=obj -Wno-gnu-line-marker) add_example_executable(example_grouped_conv_bwd_weight_xdl_bf16 grouped_conv_bwd_weight_xdl_bf16.cpp) add_example_dependencies(example_grouped_conv_bwd_weight example_grouped_conv_bwd_weight_xdl_bf16) diff --git a/example/20_grouped_conv_bwd_weight/common.hpp b/example/20_grouped_conv_bwd_weight/common.hpp index e0034bf7eb..2386e35b48 100644 --- a/example/20_grouped_conv_bwd_weight/common.hpp +++ b/example/20_grouped_conv_bwd_weight/common.hpp @@ -41,13 +41,15 @@ struct CommonLayoutSetting using WeightLayout = WeightLay; using OutputLayout = OutputLay; }; - +using ALayout = ck::tensor_layout::convolution::NHWGC; +using BLayout = ck::tensor_layout::convolution::GKYXC; +using ELayout = ck::tensor_layout::convolution::NHWGK; namespace ctl = ck::tensor_layout::convolution; template struct CommonLayoutSettingSelector : CommonLayoutSetting>, ck::tuple_element_t>, ck::tuple_element_t>> { }; @@ -123,7 +125,7 @@ inline bool parse_cmd_args(int argc, const ck::index_t num_dim_spatial = std::stoi(argv[4]); conv_param = ck::utils::conv::parse_conv_param( - num_dim_spatial, threshold_to_catch_partial_args, argv); + num_dim_spatial, threshold_to_catch_partial_args + 1, argv); } else { diff --git a/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp b/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp index 5c1c97d615..71ec0d75d6 100644 --- a/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp +++ b/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp @@ -3,7 +3,9 @@ #include "common.hpp" -#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp" +//#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp" +#include "ck/utility/blkgemmpipe_scheduler.hpp" +#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp" using InDataType = F16; using WeiDataType = F16; @@ -14,6 +16,7 @@ using InElementOp = PassThrough; using WeiElementOp = PassThrough; using OutElementOp = PassThrough; +#if 0 template using DeviceConvBwdWeightInstance = ck::tensor_operation::device::DeviceGroupedConvBwdWeight_Xdl_CShuffle< @@ -51,20 +54,86 @@ using DeviceConvBwdWeightInstance = S<0, 3, 1, 2>, // ABlockTransferThreadClusterArrangeOrder S<0, 2, 1, 3>, // ABlockTransferSrcAccessOrder 2, // ABlockTransferSrcVectorDim - 8, // ABlockTransferSrcScalarPerVector - 2, // ABlockTransferDstScalarPerVector_K1 + 1, // ABlockTransferSrcScalarPerVector + 1, // ABlockTransferDstScalarPerVector_K1 true, // ABlockLdsAddExtraM S<1, 4, 16, 4>, // BBlockTransferThreadClusterLengths_K0_N_K1 S<0, 3, 1, 2>, // BBlockTransferThreadClusterArrangeOrder S<0, 2, 1, 3>, // BBlockTransferSrcAccessOrder 2, // BBlockTransferSrcVectorDim - 8, // BBlockTransferSrcScalarPerVector - 2, // BBlockTransferDstScalarPerVector_K1 + 1, // BBlockTransferSrcScalarPerVector + 1, // BBlockTransferDstScalarPerVector_K1 true, // BBlockLdsAddExtraN 1, // CShuffleMXdlPerWavePerShuffle 1, // CShuffleNXdlPerWavePerShuffle S<1, 32, 1, 4>, // CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock - 128 / (sizeof(WeiDataType) * CHAR_BIT)>; // CBlockTransferScalarPerVector_NWaveNPerXdl + 1>; // CBlockTransferScalarPerVector_NWaveNPerXdl + +#endif + +// DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, F32, PassThrough, PassThrough, PassThrough, ConvSpec, 64, 32, 32, 32, 8, 32, 32, 1, 1, S<4, 8, 1>, S<2, 0, 1>, S<1, 0, 2>, 1, 2, 2, false, S<4, 16, 1>, S<2, 0, 1>, S<1, 0, 2>, 1, 2, 2, false, 1, 1, S<1, 8, 1, 8>, 1, Scheduler, PipelineVersion, 2>, + + //ConvBwdWeightDefault, +//is_NHWGC_GKYXC_NHWGK +using ALayout = ck::tensor_layout::convolution::NHWGC; +using BLayout = ck::tensor_layout::convolution::GKYXC; +using ELayout = ck::tensor_layout::convolution::NHWGK; +//using Scheduler =ck::BlockGemmPipelineScheduler::Intrawave; +//using PipelineVersion =ck::BlockGemmPipelineVersion::v1; +template +using DeviceConvBwdWeightInstance = + ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle< + NDimSpatial, + ALayout, + BLayout, + ELayout, + F16, + F16, + F16, + F32, + PassThrough, + PassThrough, + PassThrough, + ConvBwdWeightDefault, + 64, + 32, + 64, + 32, + 8, + 32, + 32, + 1, + 2, + S<4, 8, 1>, + S<2, 0, 1>, + S<1, 0, 2>, + 1, + 2, + 2, + false, + S<4, 16, 1>, + S<2, 0, 1>, + S<1, 0, 2>, + 1, + 2, + 2, + false, + 1, + 1, + S<1, 8, 1, 8>, + 1, + ck::BlockGemmPipelineScheduler::Intrawave, + ck::BlockGemmPipelineVersion::v2, + 2>; +#if 0 + + DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle< NDimSpatial, ALayout, BLayout, ELayout, F16, F16, F16, F32, PassThrough, PassThrough, PassThrough, ConvSpec, 64, 32, 64, 32, 8, 32, 32, 1, 2, S<4, 8, 1>, S<2, 0, 1>, S<1, 0, 2>, 1, 2, 2, false, S<4, 16, 1>, S<2, 0, 1>, S<1, 0, 2>, 1, 2, 2, false, 1, 1, S<1, 8, 1, 8>, 1, Scheduler, PipelineVersion, 2>, + + +64, 16, 16, 32, 8, 16, 16, 1, 1, S<4, 8, 1>, S<2, 0, 1>, S<1, 0, 2>, 1, 1, 4, false, S<4, 8, 1>, + S<2, 0, 1>, S<1, 0, 2>, 1, 1, 4, false, 1, 1, S<1, 8, 1, 8>, 1, + ck::BlockGemmPipelineScheduler::Intrawave, ck::BlockGemmPipelineVersion::v1, 8 > ; +#endif template using HostConvBwdWeightInstance = ck::tensor_operation::host::ReferenceConvBwdWeight(config, conv_param); + case 1: break;//return !run_grouped_conv_bwd_weight<1>(config, conv_param); case 2: return !run_grouped_conv_bwd_weight<2>(config, conv_param); - case 3: return !run_grouped_conv_bwd_weight<3>(config, conv_param); + case 3: break;//return !run_grouped_conv_bwd_weight<3>(config, conv_param); default: break; } diff --git a/example/20_grouped_conv_bwd_weight/run_grouped_conv_bwd_weight_example.inc b/example/20_grouped_conv_bwd_weight/run_grouped_conv_bwd_weight_example.inc index f320c0305b..3780c3559e 100644 --- a/example/20_grouped_conv_bwd_weight/run_grouped_conv_bwd_weight_example.inc +++ b/example/20_grouped_conv_bwd_weight/run_grouped_conv_bwd_weight_example.inc @@ -95,6 +95,9 @@ bool run_grouped_conv_bwd_weight(const ExecutionConfig& config, WeiElementOp{}, OutElementOp{}, split_k); + invoker.ShowInfo(argument); + DeviceMem gemm_workspace_dev(conv.GetWorkSpaceSize(&argument)); + conv.SetWorkSpacePointer(&argument, gemm_workspace_dev.GetDeviceBuffer()); if(!conv.IsSupportedArgument(argument)) { @@ -104,7 +107,7 @@ bool run_grouped_conv_bwd_weight(const ExecutionConfig& config, return true; } - invoker.Run(argument, StreamConfig{nullptr, false}); + invoker.Run(argument, StreamConfig{nullptr, config.time_kernel}); if(config.do_verification) { diff --git a/include/ck/host_utility/device_prop.hpp b/include/ck/host_utility/device_prop.hpp index 5439bbe1f0..b18dd29195 100644 --- a/include/ck/host_utility/device_prop.hpp +++ b/include/ck/host_utility/device_prop.hpp @@ -54,6 +54,7 @@ inline std::string get_device_name() inline bool is_xdl_supported() { + return true; return ck::get_device_name() == "gfx908" || ck::get_device_name() == "gfx90a" || ck::get_device_name() == "gfx942" || ck::get_device_name() == "gfx950"; } diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp index c7d95254c5..a8d249a5a3 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp @@ -1422,6 +1422,8 @@ struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle float Run(const Argument& arg, const StreamConfig& stream_config = StreamConfig{}) { +#if 1 +//#ifdef NDEBUG float avg_time = 0.f; auto launch_elementwise_kernel = [&]() { const AccDataType* p_c_grid = type_convert(arg.p_workspace_); @@ -1543,6 +1545,11 @@ struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle avg_time += RunGemmV3(arg, stream_config); avg_time += launch_elementwise_kernel(); return avg_time; +#else + ignore = arg; + ignore = stream_config; + return 0; +#endif } float Run(const BaseArgument* p_arg, @@ -1596,6 +1603,8 @@ struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle } if constexpr(NDimSpatial == 2) { + static_assert(is_NHWGC_GKYXC_NHWGK() || + is_NGCHW_NGKHW()); if constexpr(!(is_NHWGC_GKYXC_NHWGK() || is_NGCHW_NGKHW())) { @@ -1632,8 +1641,10 @@ struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle if constexpr(NumGroupsToMerge > 1) { // support only if whole M and N can be proccessed on one block + //static_assert(GemmM <= MPerBlock && GemmN <= NPerBlock); if(!(GemmM <= MPerBlock && GemmN <= NPerBlock)) { + printf("%d, %d, %d, %d\n",GemmM,MPerBlock, GemmN,NPerBlock); return false; } if(!(arg.Conv_C_ == 1 && arg.Conv_K_ == 1)) @@ -1669,6 +1680,7 @@ struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle } // vector load A/B matrix from global memory + static_assert(ABlockTransferSrcVectorDim == 1 && BBlockTransferSrcVectorDim == 1); if(!(ABlockTransferSrcVectorDim == 1 && BBlockTransferSrcVectorDim == 1)) { return false; diff --git a/include/ck/utility/generic_memory_space_atomic.hpp b/include/ck/utility/generic_memory_space_atomic.hpp index ab9cc4199c..e483813bf3 100644 --- a/include/ck/utility/generic_memory_space_atomic.hpp +++ b/include/ck/utility/generic_memory_space_atomic.hpp @@ -12,7 +12,7 @@ namespace ck { // instantiate this template. The purpose is to make the implementation of atomic_add explicit for // each datatype. template -__device__ X atomic_add(X* p_dst, const X& x); +__device__ X atomic_add(X* p_dst, const X& x) { ignore = p_dst; ignore = x; return 0;}; template <> __device__ int32_t atomic_add(int32_t* p_dst, const int32_t& x)