diff --git a/example/01_gemm/CMakeLists.txt b/example/01_gemm/CMakeLists.txt index 96678d275a..37a583b6e7 100755 --- a/example/01_gemm/CMakeLists.txt +++ b/example/01_gemm/CMakeLists.txt @@ -18,13 +18,17 @@ endif(USE_BITINT_EXTENSION_INT4) add_custom_target(example_gemm_xdl) add_example_executable(example_gemm_xdl_fp16 gemm_xdl_fp16.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16) +target_compile_options(example_gemm_xdl_fp16 PRIVATE -save-temps=obj -Wno-gnu-line-marker) add_example_executable(example_gemm_xdl_fp16_v2 gemm_xdl_fp16_v2.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16_v2) +target_compile_options(example_gemm_xdl_fp16_v2 PRIVATE -save-temps=obj -Wno-gnu-line-marker) add_example_executable(example_gemm_xdl_fp16_streamk_v3 gemm_xdl_fp16_streamk_v3.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16_streamk_v3) add_example_executable(example_gemm_xdl_fp16_v3 gemm_xdl_fp16_v3.cpp) +target_compile_options(example_gemm_xdl_fp16_v3 PRIVATE -save-temps=obj -Wno-gnu-line-marker) + add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16_v3) add_example_executable(example_gemm_xdl_fp8_v3 gemm_xdl_fp8_v3.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp8_v3) @@ -38,6 +42,7 @@ add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16_fp8_streamk_v3) add_example_executable(example_gemm_xdl_bf16_v3 gemm_xdl_bf16_v3.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_bf16_v3) +target_compile_options(example_gemm_xdl_bf16_v3 PRIVATE -save-temps=obj -Wno-gnu-line-marker) list(APPEND gpu_list gfx942 gfx950) set(target 0) @@ -60,6 +65,7 @@ add_example_dependencies(example_gemm_xdl example_gemm_xdl_skip_b_lds_fp16) add_example_executable(example_gemm_xdl_bf16 gemm_xdl_bf16.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_bf16) +target_compile_options(example_gemm_xdl_bf16 PRIVATE -save-temps=obj -Wno-gnu-line-marker) add_example_executable(example_gemm_xdl_int8 gemm_xdl_int8.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_int8) diff --git a/example/01_gemm/gemm_xdl_fp16_v2.cpp b/example/01_gemm/gemm_xdl_fp16_v2.cpp index ecd3b7be5d..c71ec7ad4b 100644 --- a/example/01_gemm/gemm_xdl_fp16_v2.cpp +++ b/example/01_gemm/gemm_xdl_fp16_v2.cpp @@ -59,4 +59,80 @@ using ReferenceGemmInstanceGPU = ck::tensor_operation::device::ReferenceGemm, S<1, 0, 2>, S<1, 0, 2>, + 2, 8, 8, 0, + S<8, 32, 1>, S<0, 2, 1>, S<0, 2, 1>, + 1, 8, 4, 0, + 1, 1, S<1, 32, 1, 8>, 8, + ck::LoopScheduler::Default, ck::PipelineVersion::v1>; + +using DeviceGemmInstance_1 = + ck::tensor_operation::device::DeviceGemm_Xdl_CShuffleV2< + ALayout, BLayout, CLayout, + F16, F16, F16, F32, F16, + PassThrough, PassThrough, PassThrough, GemmDefault, + 2, 256, + 256, 256, + 32, 8, 8, + 32, 32, + 4, 4, + S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, + 2, 8, 8, 0, + S<4, 32, 2>, S<0, 2, 1>, S<0, 2, 1>, + 1, 8, 8, 0, + 1, 1, S<1, 32, 1, 8>, 8, + ck::LoopScheduler::Default, ck::PipelineVersion::v1>; + +using DeviceGemmInstance_2 = + ck::tensor_operation::device::DeviceGemm_Xdl_CShuffleV2< + ALayout, BLayout, CLayout, + F16, F16, F16, F32, F16, + PassThrough, PassThrough, PassThrough, GemmDefault, + 2, 256, + 256, 256, + 32, 8, 8, + 32, 32, + 4, 4, + S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, + 2, 8, 8, 0, + S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, + 2, 8, 8, 0, + 1, 1, S<1, 32, 1, 8>, 8, + ck::LoopScheduler::Default, ck::PipelineVersion::v1>; + +int main(int argc, char* argv[]) +{ + run_gemm_example_with_instance(argc, argv); + run_gemm_example_with_instance(argc, argv); + run_gemm_example_with_instance(argc, argv); +} diff --git a/example/01_gemm/gemm_xdl_fp16_v3.cpp b/example/01_gemm/gemm_xdl_fp16_v3.cpp index 4a969246cd..a8777d3fc7 100644 --- a/example/01_gemm/gemm_xdl_fp16_v3.cpp +++ b/example/01_gemm/gemm_xdl_fp16_v3.cpp @@ -19,7 +19,7 @@ using AElementOp = PassThrough; using BElementOp = PassThrough; using CElementOp = PassThrough; -static constexpr auto GemmDefault = ck::tensor_operation::device::GemmSpecialization::MNPadding; +static constexpr auto GemmDefault = ck::tensor_operation::device::GemmSpecialization::Default; // clang-format off using DeviceGemmV2Instance = @@ -27,17 +27,17 @@ using DeviceGemmV2Instance = ALayout, BLayout, CLayout, ADataType, BDataType, CDataType, AccDataType, CShuffleDataType, PassThrough, PassThrough, PassThrough, GemmDefault, - 64, - 16, 16, - 256, 8, 8, + 256, + 128, 128, + 64, 8, 8, 16, 16, - 1, 1, - S<32, 2, 1>, S<1, 0, 2>, S<1, 0, 2>, + 4, 4, + S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 0, - S<32, 2, 1>, S<1, 0, 2>, S<1, 0, 2>, + S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 0, - 1, 1, S<1, 16, 1, 4>, 4, - ck::BlockGemmPipelineScheduler::Interwave, ck::BlockGemmPipelineVersion::v2>; + 1, 2, S<1, 32, 1, 8>, 8, + ck::BlockGemmPipelineScheduler::Intrawave,ck::BlockGemmPipelineVersion::v3>; // clang-format on using ReferenceGemmInstance = ck::tensor_operation::host:: diff --git a/example/01_gemm/run_gemm_example.inc b/example/01_gemm/run_gemm_example.inc index c064ed500c..9a70174e9a 100644 --- a/example/01_gemm/run_gemm_example.inc +++ b/example/01_gemm/run_gemm_example.inc @@ -3,7 +3,7 @@ #pragma once -template +template bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) { #if defined(BUILD_INT4_EXAMPLE) && defined(CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4) @@ -19,6 +19,13 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) auto StrideB = problem_size.StrideB; auto StrideC = problem_size.StrideC; + using ALayout_ = typename GemmConfig::ALayout_; + using BLayout_ = typename GemmConfig::BLayout_; + using CLayout_ = typename GemmConfig::CLayout_; + using ADataType_ = typename GemmConfig::ADataType_; + using BDataType_ = typename GemmConfig::BDataType_; + using CDataType_ = typename GemmConfig::CDataType_; + auto f_host_tensor_descriptor = [](std::size_t row, std::size_t col, std::size_t stride, auto layout) { if constexpr(std::is_same_v) @@ -49,68 +56,68 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) return static_cast(stride); }; - StrideA = f_get_default_stride(M, K, StrideA, ALayout{}); - StrideB = f_get_default_stride(K, N, StrideB, BLayout{}); - StrideC = f_get_default_stride(M, N, StrideC, CLayout{}); + StrideA = f_get_default_stride(M, K, StrideA, ALayout_{}); + StrideB = f_get_default_stride(K, N, StrideB, BLayout_{}); + StrideC = f_get_default_stride(M, N, StrideC, CLayout_{}); - Tensor a_m_k(f_host_tensor_descriptor(M, K, StrideA, ALayout{})); - Tensor b_k_n(f_host_tensor_descriptor(K, N, StrideB, BLayout{})); + Tensor a_m_k(f_host_tensor_descriptor(M, K, StrideA, ALayout_{})); + Tensor b_k_n(f_host_tensor_descriptor(K, N, StrideB, BLayout_{})); switch(config.init_method) { case 0: - ck::utils::FillConstant{ck::type_convert(1.f)}(a_m_k); - ck::utils::FillConstant{ck::type_convert(1.f)}(b_k_n); + ck::utils::FillConstant{ck::type_convert(1.f)}(a_m_k); + ck::utils::FillConstant{ck::type_convert(1.f)}(b_k_n); break; case 1: - ck::utils::FillUniformDistributionIntegerValue{-5.f, 5.f}(a_m_k); - ck::utils::FillUniformDistributionIntegerValue{-5.f, 5.f}(b_k_n); + ck::utils::FillUniformDistributionIntegerValue{-5.f, 5.f}(a_m_k); + ck::utils::FillUniformDistributionIntegerValue{-5.f, 5.f}(b_k_n); break; case 2: - ck::utils::FillUniformDistribution{-1.f, 1.f}(a_m_k); - ck::utils::FillUniformDistribution{-1.f, 1.f}(b_k_n); + ck::utils::FillUniformDistribution{-1.f, 1.f}(a_m_k); + ck::utils::FillUniformDistribution{-1.f, 1.f}(b_k_n); break; case 3: - ck::utils::FillUniformDistributionIntegerValue{1.f, 1.f}(a_m_k); - ck::utils::FillUniformDistributionIntegerValue{-5.f, 5.f}(b_k_n); + ck::utils::FillUniformDistributionIntegerValue{1.f, 1.f}(a_m_k); + ck::utils::FillUniformDistributionIntegerValue{-5.f, 5.f}(b_k_n); break; case 4: - ck::utils::FillUniformDistributionIntegerValue{-5.f, 5.f}(a_m_k); - ck::utils::FillUniformDistributionIntegerValue{1.f, 1.f}(b_k_n); + ck::utils::FillUniformDistributionIntegerValue{-5.f, 5.f}(a_m_k); + ck::utils::FillUniformDistributionIntegerValue{1.f, 1.f}(b_k_n); break; case 5: - ck::utils::FillUniformDistributionIntegerValue{-2.f, 2.f}(a_m_k); - ck::utils::FillUniformDistributionIntegerValue{-2.f, 2.f}(b_k_n); + ck::utils::FillUniformDistributionIntegerValue{-2.f, 2.f}(a_m_k); + ck::utils::FillUniformDistributionIntegerValue{-2.f, 2.f}(b_k_n); break; default: - ck::utils::FillUniformDistribution{-0.1f, 0.1f}(a_m_k); - ck::utils::FillUniformDistribution{-0.1f, 0.1f}(b_k_n); + ck::utils::FillUniformDistribution{-0.1f, 0.1f}(a_m_k); + ck::utils::FillUniformDistribution{-0.1f, 0.1f}(b_k_n); } - Tensor c_m_n_host_result(f_host_tensor_descriptor(M, N, StrideC, CLayout{})); - Tensor c_m_n_device_result(f_host_tensor_descriptor(M, N, StrideC, CLayout{})); - Tensor c_m_n_device_ref_result(f_host_tensor_descriptor(M, N, StrideC, CLayout{})); + Tensor c_m_n_host_result(f_host_tensor_descriptor(M, N, StrideC, CLayout_{})); + Tensor c_m_n_device_result(f_host_tensor_descriptor(M, N, StrideC, CLayout_{})); + Tensor c_m_n_device_ref_result(f_host_tensor_descriptor(M, N, StrideC, CLayout_{})); std::cout << "a_m_k: " << a_m_k.mDesc << std::endl; std::cout << "b_k_n: " << b_k_n.mDesc << std::endl; std::cout << "c_m_n: " << c_m_n_host_result.mDesc << std::endl; #ifdef BUILD_INT4_EXAMPLE - DeviceMem a_m_k_device_buf(sizeof(KernelADataType) * a_m_k.mDesc.GetElementSpaceSize()); + DeviceMem a_m_k_device_buf(sizeof(KernelADataType_) * a_m_k.mDesc.GetElementSpaceSize()); DeviceMem b_k_n_device_buf(sizeof(KernelBDataType) * b_k_n.mDesc.GetElementSpaceSize()); DeviceMem c_m_n_device_buf(sizeof(KernelCDataType) * c_m_n_device_result.mDesc.GetElementSpaceSize()); - const Tensor a_m_k_converted(a_m_k); + const Tensor a_m_k_converted(a_m_k); const Tensor b_k_n_converted(b_k_n); a_m_k_device_buf.ToDevice(a_m_k_converted.mData.data()); b_k_n_device_buf.ToDevice(b_k_n_converted.mData.data()); #else - DeviceMem a_m_k_device_buf(sizeof(ADataType) * a_m_k.mDesc.GetElementSpaceSize()); - DeviceMem b_k_n_device_buf(sizeof(BDataType) * b_k_n.mDesc.GetElementSpaceSize()); - DeviceMem c_m_n_device_buf(sizeof(CDataType) * c_m_n_device_result.mDesc.GetElementSpaceSize()); - DeviceMem c_m_n_device_ref_buf(sizeof(CDataType) * + DeviceMem a_m_k_device_buf(sizeof(ADataType_) * a_m_k.mDesc.GetElementSpaceSize()); + DeviceMem b_k_n_device_buf(sizeof(BDataType_) * b_k_n.mDesc.GetElementSpaceSize()); + DeviceMem c_m_n_device_buf(sizeof(CDataType_) * c_m_n_device_result.mDesc.GetElementSpaceSize()); + DeviceMem c_m_n_device_ref_buf(sizeof(CDataType_) * c_m_n_device_ref_result.mDesc.GetElementSpaceSize()); a_m_k_device_buf.ToDevice(a_m_k.mData.data()); @@ -123,7 +130,7 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) auto c_element_op = CElementOp{}; // do GEMM - auto gemm = DeviceGemmInstance{}; + auto gemm = DeviceGemm{}; auto invoker = gemm.MakeInvoker(); float ave_time = 0; @@ -131,13 +138,13 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) { auto argument = gemm.MakeArgument( #ifdef BUILD_INT4_EXAMPLE - static_cast(a_m_k_device_buf.GetDeviceBuffer()), + static_cast(a_m_k_device_buf.GetDeviceBuffer()), static_cast(b_k_n_device_buf.GetDeviceBuffer()), static_cast(c_m_n_device_buf.GetDeviceBuffer()), #else - static_cast(a_m_k_device_buf.GetDeviceBuffer()), - static_cast(b_k_n_device_buf.GetDeviceBuffer()), - static_cast(c_m_n_device_buf.GetDeviceBuffer()), + static_cast(a_m_k_device_buf.GetDeviceBuffer()), + static_cast(b_k_n_device_buf.GetDeviceBuffer()), + static_cast(c_m_n_device_buf.GetDeviceBuffer()), #endif M, N, @@ -169,7 +176,7 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) std::size_t flop = 2_uz * M * N * K; std::size_t num_btype = - sizeof(ADataType) * M * K + sizeof(BDataType) * K * N + sizeof(CDataType) * M * N; + sizeof(ADataType_) * M * K + sizeof(BDataType_) * K * N + sizeof(CDataType_) * M * N; float tflops = static_cast(flop) / 1.E9 / ave_time; @@ -193,11 +200,11 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) ref_invoker.Run(ref_argument); #ifdef BUILD_INT4_EXAMPLE - Tensor c_m_n_device_result_converted(c_m_n_host_result.mDesc); + Tensor c_m_n_device_result_converted(c_m_n_host_result.mDesc); c_m_n_device_buf.FromDevice(c_m_n_device_result_converted.mData.data()); - c_m_n_device_result = c_m_n_device_result_converted.CopyAsType(); + c_m_n_device_result = c_m_n_device_result_converted.CopyAsType(); return ck::utils::check_err(c_m_n_device_result_converted, c_m_n_host_result); #else @@ -206,8 +213,8 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) pass &= ck::utils::check_err(c_m_n_device_result, c_m_n_host_result, "Error: Incorrect results!", - get_rtol(), - get_atol()); + get_rtol(), + get_atol()); #endif } @@ -218,9 +225,9 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) auto ref_invoker_gpu = ref_gemm_gpu.MakeInvoker(); auto ref_argument_gpu = ref_gemm_gpu.MakeArgument( - static_cast(a_m_k_device_buf.GetDeviceBuffer()), - static_cast(b_k_n_device_buf.GetDeviceBuffer()), - static_cast(c_m_n_device_ref_buf.GetDeviceBuffer()), + static_cast(a_m_k_device_buf.GetDeviceBuffer()), + static_cast(b_k_n_device_buf.GetDeviceBuffer()), + static_cast(c_m_n_device_ref_buf.GetDeviceBuffer()), M, N, K, @@ -237,17 +244,36 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) pass &= ck::utils::check_err(c_m_n_device_result, c_m_n_device_ref_result, "Error: Incorrect results!", - get_rtol(), - get_atol()); + get_rtol(), + get_atol()); } return pass == true; } +struct GemmConfigDefault +{ + using ALayout_ = ALayout; + using BLayout_ = BLayout; + using CLayout_ = CLayout; + using ADataType_ = ADataType; + using BDataType_ = BDataType; + using CDataType_ = CDataType; +}; + bool run_gemm_example(int argc, char* argv[]) { ProblemSize problem_size; ExecutionConfig config; - return !parse_cmd_args(argc, argv, problem_size, config) || run_gemm(problem_size, config); + return !parse_cmd_args(argc, argv, problem_size, config) || run_gemm(problem_size, config); +} + +template +bool run_gemm_example_with_instance(int argc, char* argv[]) +{ + ProblemSize problem_size; + ExecutionConfig config; + + return !parse_cmd_args(argc, argv, problem_size, config) || run_gemm(problem_size, config); } diff --git a/example/ck_tile/03_gemm/CMakeLists.txt b/example/ck_tile/03_gemm/CMakeLists.txt index 411db2e317..935a29556b 100644 --- a/example/ck_tile/03_gemm/CMakeLists.txt +++ b/example/ck_tile/03_gemm/CMakeLists.txt @@ -7,3 +7,8 @@ endif() list(APPEND EXAMPLE_GEMM_COMPILE_OPTIONS -mllvm -enable-noalias-to-md-conversion=0) target_compile_options(tile_example_gemm_basic PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS}) target_compile_options(tile_example_gemm_universal PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS}) + +target_compile_options(tile_example_gemm_universal PRIVATE -save-temps=obj -Wno-gnu-line-marker) +target_compile_options(tile_example_gemm_basic PRIVATE -save-temps=obj -Wno-gnu-line-marker) +#target_compile_options(tile_example_gemm_basic PRIVATE ";-mllvm;--slp-threshold=-32") +#target_compile_options(tile_example_gemm_universal PRIVATE ";-mllvm;--slp-threshold=-32") \ No newline at end of file diff --git a/example/ck_tile/03_gemm/gemm_basic.cpp b/example/ck_tile/03_gemm/gemm_basic.cpp index 1edb3da947..e50c83e895 100644 --- a/example/ck_tile/03_gemm/gemm_basic.cpp +++ b/example/ck_tile/03_gemm/gemm_basic.cpp @@ -26,11 +26,11 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& constexpr bool kPadN = false; constexpr bool kPadK = false; - constexpr int kBlockPerCu = 1; + constexpr int kBlockPerCu = 2; // This part comes from the Codegen - constexpr ck_tile::index_t M_Tile = 256; - constexpr ck_tile::index_t N_Tile = 256; + constexpr ck_tile::index_t M_Tile = 128; + constexpr ck_tile::index_t N_Tile = 128; constexpr ck_tile::index_t K_Tile = 64; constexpr ck_tile::index_t M_Warp = 2; @@ -144,16 +144,17 @@ int run_gemm_example_prec_type(std::string a_layout, std::string b_layout, int a } else { - if(a_layout == "R" && b_layout == "R") - { - return run_gemm_example_with_layouts( - argc, argv, Row{}, Row{}, Row{}); - } - else if(a_layout == "R" && b_layout == "C") + if(a_layout == "R" && b_layout == "C") { return run_gemm_example_with_layouts( argc, argv, Row{}, Col{}, Row{}); } +#if 0 + else if(a_layout == "R" && b_layout == "R") + { + return run_gemm_example_with_layouts( + argc, argv, Row{}, Row{}, Row{}); + } else if(a_layout == "C" && b_layout == "R") { return run_gemm_example_with_layouts( @@ -164,6 +165,7 @@ int run_gemm_example_prec_type(std::string a_layout, std::string b_layout, int a return run_gemm_example_with_layouts( argc, argv, Col{}, Col{}, Row{}); } +#endif else { throw std::runtime_error("Unsupported memory layout for the input matrices!"); @@ -185,6 +187,7 @@ int run_gemm_example(int argc, char* argv[]) { return run_gemm_example_prec_type(a_layout, b_layout, argc, argv); } +#if 0 else if(data_type == "bf16") { return run_gemm_example_prec_type(a_layout, b_layout, argc, argv); @@ -207,6 +210,7 @@ int run_gemm_example(int argc, char* argv[]) return run_gemm_example_prec_type( a_layout, b_layout, argc, argv); } +#endif #endif else { diff --git a/example/ck_tile/03_gemm/gemm_utils.hpp b/example/ck_tile/03_gemm/gemm_utils.hpp index 25fab6bde0..45c2513733 100644 --- a/example/ck_tile/03_gemm/gemm_utils.hpp +++ b/example/ck_tile/03_gemm/gemm_utils.hpp @@ -57,7 +57,7 @@ struct GemmConfig // Compute friendly for Intrawave scheduler static constexpr ck_tile::index_t M_Tile = 128; static constexpr ck_tile::index_t N_Tile = 128; - static constexpr ck_tile::index_t K_Tile = 128; + static constexpr ck_tile::index_t K_Tile = 64; static constexpr ck_tile::index_t M_Warp = 2; static constexpr ck_tile::index_t N_Warp = 2; @@ -96,7 +96,7 @@ struct GemmConfig static constexpr bool TransposeC = false; static constexpr bool UseStructuredSparsity = false; - static constexpr int kBlockPerCu = 1; + static constexpr int kBlockPerCu = 2; static constexpr ck_tile::index_t TileParitionerGroupNum = 8; static constexpr ck_tile::index_t TileParitionerM01 = 4; }; diff --git a/example/ck_tile/03_gemm/universal_gemm.cpp b/example/ck_tile/03_gemm/universal_gemm.cpp index b60a3b274b..6c1e322623 100644 --- a/example/ck_tile/03_gemm/universal_gemm.cpp +++ b/example/ck_tile/03_gemm/universal_gemm.cpp @@ -41,10 +41,9 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& sequence, GemmConfig::PermuteA, GemmConfig::PermuteB>; - using TilePartitioner = - ck_tile::GemmSpatiallyLocalTilePartitioner; + using TilePartitioner = ck_tile::GemmSpatiallyLocalTilePartitioner; using Traits = ck_tile::TileGemmTraits{}); } + #if 0 else { Run(has_hot_loop_, @@ -148,6 +148,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& ck_tile::integral_constant{}); } + #endif }; if(has_hot_loop) @@ -158,6 +159,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& RunSplitk(ck_tile::bool_constant{}, ck_tile::integral_constant{}); } + #if 0 else if(tail_num == ck_tile::TailNumber::Odd) { RunSplitk(ck_tile::bool_constant{}, @@ -168,6 +170,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& RunSplitk(ck_tile::bool_constant{}, ck_tile::integral_constant{}); } + #endif else { std::ostringstream err; @@ -215,6 +218,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& } else { + #if 0 if(tail_num == ck_tile::TailNumber::Full) { RunSplitk(ck_tile::bool_constant{}, @@ -231,6 +235,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& ck_tile::integral_constant{}); } else + #endif { std::ostringstream err; err << "Num K loop must be larger than number of prefetech stages." @@ -271,16 +276,18 @@ int run_gemm_example_prec_type(std::string a_layout, std::string b_layout, int a } else { - if(a_layout == "R" && b_layout == "R") - { - return run_gemm_example_with_layouts( - argc, argv, Row{}, Row{}, Row{}); - } - else if(a_layout == "R" && b_layout == "C") + if(a_layout == "R" && b_layout == "C") { return run_gemm_example_with_layouts( argc, argv, Row{}, Col{}, Row{}); } +#if 0 + else if(a_layout == "R" && b_layout == "R") + { + return run_gemm_example_with_layouts( + argc, argv, Row{}, Row{}, Row{}); + } + else if(a_layout == "C" && b_layout == "R") { return run_gemm_example_with_layouts( @@ -291,6 +298,7 @@ int run_gemm_example_prec_type(std::string a_layout, std::string b_layout, int a return run_gemm_example_with_layouts( argc, argv, Col{}, Col{}, Row{}); } +#endif else { throw std::runtime_error("Unsupported memory layout for the input matrices!"); @@ -312,6 +320,7 @@ int run_gemm_example(int argc, char* argv[]) { return run_gemm_example_prec_type(a_layout, b_layout, argc, argv); } +#if 0 else if(data_type == "bf16") { return run_gemm_example_prec_type(a_layout, b_layout, argc, argv); @@ -334,6 +343,7 @@ int run_gemm_example(int argc, char* argv[]) return run_gemm_example_prec_type( a_layout, b_layout, argc, argv); } +#endif #endif else { diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops.hpp index 438d7d8ac3..0ad5b205dd 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops.hpp @@ -9,6 +9,7 @@ #include "ck/tensor_operation/gpu/warp/xdlops_gemm.hpp" #include "ck/tensor_description/tensor_adaptor.hpp" +#define ENABLE_HOT 1 // Double LDS buffer // Prefetech 2 stage // Local prefetch 1 stage @@ -372,6 +373,7 @@ struct BlockwiseGemmXdlops_pipeline_v4 __device__ static constexpr auto HotLoopScheduler() { + #if ENABLE_HOT // schedule constexpr auto num_ds_read_inst = HotLoopInstList::A_LDS_Read_Inst_Num + HotLoopInstList::B_LDS_Read_Inst_Num; @@ -398,13 +400,14 @@ struct BlockwiseGemmXdlops_pipeline_v4 __builtin_amdgcn_sched_group_barrier( 0x008, num_mfma_inst / num_buffer_load_inst - 3, 0); // MFMA }); + #endif } template __device__ static constexpr auto TailScheduler() { } - +#if ENABLE_HOT template <> __device__ constexpr auto TailScheduler<1>() { @@ -449,6 +452,7 @@ struct BlockwiseGemmXdlops_pipeline_v4 0x008, num_mfma_inst / num_ds_read_inst, 0); // MFMA }); } +#endif static constexpr AMmaTileDesc a_block_desc_m0_m1_m2_k; static constexpr BMmaTileDesc b_block_desc_n0_n1_n2_k; diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp index dde21725d0..17ef6a929e 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp @@ -357,6 +357,7 @@ struct DeviceGemm_Xdl_CShuffleV3 : public DeviceGemmV2 1) { const auto kernel = @@ -367,6 +368,7 @@ struct DeviceGemm_Xdl_CShuffleV3 : public DeviceGemmV2