From 349f6c073f01de47327f84f993ed91683eb901ee Mon Sep 17 00:00:00 2001 From: coderfeli Date: Tue, 4 Mar 2025 01:51:31 +0000 Subject: [PATCH] checkout develop int4 related files --- CMakeLists.txt | 12 + cmake/EnableCompilerWarnings.cmake | 1 + .../problem.hpp | 35 +- ...mm_softmax_gemm_operation_xdl_cshuffle.cpp | 14 +- codegen/test/batched_gemm_softmax_gemm.cpp | 8 +- codegen/test/rtc/src/compile_kernel.cpp | 1 + example/01_gemm/CMakeLists.txt | 2 - example/01_gemm/common.hpp | 23 - .../gemm_xdl_fp8_pk_i4_bpreshuffle_v3.cpp | 369 ------- example/01_gemm/gemm_xdl_fp8_pk_i4_v3.cpp | 348 ------- .../moe_pk_i4_gemm1.cpp | 523 ---------- .../moe_pk_i4_gemm2.cpp | 491 --------- ...ipeline_xdlops_b_preshuffle_dequant_v1.hpp | 547 ----------- ...ipeline_xdlops_b_preshuffle_dequant_v3.hpp | 928 ------------------ ..._pipeline_xdlops_b_preshuffle_selector.hpp | 54 - include/ck/utility/amd_ck_fp8.hpp | 9 +- include/ck/utility/amd_inline_asm.hpp | 55 -- include/ck/utility/data_type.hpp | 14 +- include/ck/utility/mxf4_utils.hpp | 5 +- include/ck/utility/mxf6_utils.hpp | 8 +- include/ck/utility/mxfp_utils.hpp | 6 +- include/ck/utility/type_convert.hpp | 44 +- include/ck_tile/core/config.hpp | 2 +- .../fused_moe/kernel/moe_sorting_kernel.hpp | 18 +- .../gpu/gemm_ab_scale.hpp | 88 +- .../gpu/gemm_ab_scale/CMakeLists.txt | 7 +- ...le_xdl_f8_f8_bf16_mk_nk_mn_128_128_128.hpp | 69 +- ...k_mn_128_128_128_comp_default_instance.cpp | 6 +- ..._mn_128_128_128_comp_kpadding_instance.cpp | 6 +- ...n_128_128_128_comp_mnkpadding_instance.cpp | 37 + ...mn_128_128_128_comp_mnpadding_instance.cpp | 37 + ...mn_128_128_128_mem_v1_default_instance.cpp | 8 +- ...n_128_128_128_mem_v1_kpadding_instance.cpp | 8 +- ...128_128_128_mem_v1_mnkpadding_instance.cpp | 38 + ...f8_bf16_mk_mfma_mn_p1_padding_instance.cpp | 33 - ...f8_bf16_mk_mfma_mn_p2_padding_instance.cpp | 33 - ...f8_bf16_mk_mfma_mn_p3_padding_instance.cpp | 33 - ..._f8_f16_mk_mfma_mn_p1_padding_instance.cpp | 33 - ..._f8_f16_mk_mfma_mn_p2_padding_instance.cpp | 33 - ..._f8_f16_mk_mfma_mn_p3_padding_instance.cpp | 33 - profiler/src/profile_gemm_ab_scale.cpp | 8 +- 41 files changed, 349 insertions(+), 3678 deletions(-) delete mode 100644 example/01_gemm/gemm_xdl_fp8_pk_i4_bpreshuffle_v3.cpp delete mode 100644 example/01_gemm/gemm_xdl_fp8_pk_i4_v3.cpp delete mode 100644 example/65_gemm_multiply_multiply/moe_pk_i4_gemm1.cpp delete mode 100644 example/65_gemm_multiply_multiply/moe_pk_i4_gemm2.cpp delete mode 100644 include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v1.hpp delete mode 100644 include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v3.hpp create mode 100644 library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnkpadding_instance.cpp create mode 100644 library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnpadding_instance.cpp create mode 100644 library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_mnkpadding_instance.cpp delete mode 100644 library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p1_padding_instance.cpp delete mode 100644 library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p2_padding_instance.cpp delete mode 100644 library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p3_padding_instance.cpp delete mode 100644 library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p1_padding_instance.cpp delete mode 100644 library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p2_padding_instance.cpp delete mode 100644 library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p3_padding_instance.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 8f31267b64..e90f893de0 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -246,6 +246,13 @@ if(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER 500500000) add_compile_options("SHELL: -mllvm --lsr-drop-solution=1") endif() endif() +if(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER 600140090) + check_cxx_compiler_flag("-mllvm -enable-post-misched=0" HAS_ENABLE_POST_MISCHED) + if(HAS_ENABLE_POST_MISCHED) + message("Adding the enable-post-misched=0 compiler flag") + add_compile_options("SHELL: -mllvm -enable-post-misched=0") + endif() +endif() set(check-coerce) check_cxx_compiler_flag(" -mllvm -amdgpu-coerce-illegal-types=1" check-coerce) if(NOT WIN32 AND check-coerce AND ${hip_VERSION_FLAT} GREATER 600241132) @@ -527,6 +534,11 @@ include_directories(BEFORE ${HIP_INCLUDE_DIRS} ) +SET(BUILD_DEV ON CACHE BOOL "BUILD_DEV") +if(BUILD_DEV) + add_compile_options(-Werror) + add_compile_options(-Weverything) +endif() message("CMAKE_CXX_FLAGS: ${CMAKE_CXX_FLAGS}") if("${CMAKE_CXX_COMPILER_ID}" MATCHES "Clang") diff --git a/cmake/EnableCompilerWarnings.cmake b/cmake/EnableCompilerWarnings.cmake index d5bcd6f978..fb2b38d688 100644 --- a/cmake/EnableCompilerWarnings.cmake +++ b/cmake/EnableCompilerWarnings.cmake @@ -66,6 +66,7 @@ else() -Wunreachable-code -Wunused -Wno-reserved-identifier + -Werror -Wno-option-ignored -Wsign-compare -Wno-extra-semi-stmt diff --git a/codegen/include/ck/host/device_batched_gemm_softmax_gemm/problem.hpp b/codegen/include/ck/host/device_batched_gemm_softmax_gemm/problem.hpp index 428034a3ba..8e68f6cc88 100644 --- a/codegen/include/ck/host/device_batched_gemm_softmax_gemm/problem.hpp +++ b/codegen/include/ck/host/device_batched_gemm_softmax_gemm/problem.hpp @@ -15,23 +15,24 @@ namespace device_batched_gemm_softmax_gemm { // defines the problem specification for a GEMM operation struct Problem { - std::size_t M = 0; - std::size_t N = 0; - std::size_t K = 0; - std::size_t O = 0; - bool TransA = false; - bool TransB = false; - bool TransB1 = false; - bool TransC = false; - DataType ADataType = DataType::Half; - DataType BDataType = DataType::Half; - DataType B1DataType = DataType::Half; - DataType CDataType = DataType::Half; - std::string AElementOp = PassThrough; - std::string BElementOp = PassThrough; - std::string B1ElementOp = PassThrough; - std::string CElementOp = PassThrough; - std::string AccElementOp = Scale; + std::size_t M = 0; + std::size_t N = 0; + std::size_t K = 0; + std::size_t O = 0; + bool TransA = false; + bool TransB = false; + bool TransB1 = false; + bool TransC = false; + DataType ADataType = DataType::Half; + DataType BDataType = DataType::Half; + DataType B1DataType = DataType::Half; + DataType CDataType = DataType::Half; + std::string AElementOp = PassThrough; + std::string BElementOp = PassThrough; + std::string B1ElementOp = PassThrough; + std::string CElementOp = PassThrough; + std::string AccElementOp = Scale; + bool MaskOutUpperTriangle = false; // returns the correct device op file for the operation std::string GetIncludeHeader() const; diff --git a/codegen/src/device_batched_gemm_softmax_gemm_operation_xdl_cshuffle.cpp b/codegen/src/device_batched_gemm_softmax_gemm_operation_xdl_cshuffle.cpp index b12c2e1a4a..6029ab0c7d 100644 --- a/codegen/src/device_batched_gemm_softmax_gemm_operation_xdl_cshuffle.cpp +++ b/codegen/src/device_batched_gemm_softmax_gemm_operation_xdl_cshuffle.cpp @@ -259,10 +259,7 @@ std::vector Operation_Xdl_CShuffle::CreateOperations( x.tile_desc.gemm1_n_per_block); x.update_prologue(prologue); x.update_epilogue(epilogue); - x.mask_out_upper_triangle = true; - result.push_back(x); - - x.mask_out_upper_triangle = false; + x.mask_out_upper_triangle = prob.MaskOutUpperTriangle; result.push_back(x); } return result; @@ -273,13 +270,20 @@ std::vector Operation_Xdl_CShuffle::CreateOperations( std::vector> Operation_Xdl_CShuffle::CreateOperations(const std::string& prologue, const std::string& epilogue) { + std::vector problems; + Problem prob; prob.TransA = false; prob.TransB = true; prob.TransB1 = false; prob.TransC = false; + problems.push_back(prob); - return {CreateOperations(prob, prologue, epilogue)}; + prob.MaskOutUpperTriangle = true; + problems.push_back(prob); + + return Transform(problems, + [&](const Problem& p) { return CreateOperations(p, prologue, epilogue); }); } static const char* const DeviceBatchedGemmSoftmaxGemm_Xdl_CShuffleTemplate = diff --git a/codegen/test/batched_gemm_softmax_gemm.cpp b/codegen/test/batched_gemm_softmax_gemm.cpp index 3f0b8bfe6a..0de8dbdd51 100644 --- a/codegen/test/batched_gemm_softmax_gemm.cpp +++ b/codegen/test/batched_gemm_softmax_gemm.cpp @@ -42,7 +42,7 @@ TEST_CASE(test_problem_kernel) prob.K = 1024; prob.O = 1024; prob.TransB = true; - check_all check1, check2; + check_all check; auto a = to_gpu(generate_buffer(1024 * 1024, 0)); auto b = to_gpu(generate_buffer(1024 * 1024, 1)); auto b1 = to_gpu(generate_buffer(1024 * 1024, 2)); @@ -77,10 +77,8 @@ TEST_CASE(test_problem_kernel) k.launch(nullptr, grid_size * block_size, block_size)( a.data(), b.data(), b1.data(), c.data()); - if(solution.GetTemplateParameter("MaskOutUpperTriangle")) - CHECK(report(solution, check1(rtc::from_gpu(c)))); - else - CHECK(report(solution, check2(rtc::from_gpu(c)))); + // NOTE: Solutions where MaskOutUpperTriangle is True don't produce consistent results + CHECK(report(solution, check(rtc::from_gpu(c)))); } } diff --git a/codegen/test/rtc/src/compile_kernel.cpp b/codegen/test/rtc/src/compile_kernel.cpp index a8da88be09..262e6bae46 100644 --- a/codegen/test/rtc/src/compile_kernel.cpp +++ b/codegen/test/rtc/src/compile_kernel.cpp @@ -279,6 +279,7 @@ static kernel hiprtc_compile_kernel(const std::vector& srcs, compile_o { options.flags += " -I. -O3"; options.flags += " -std=c++17"; + options.flags += " -DCK_CODE_GEN_RTC"; options.flags += " --offload-arch=" + get_device_name(); auto cos = compile_hip_src_with_hiprtc(srcs, options); if(cos.size() != 1) diff --git a/example/01_gemm/CMakeLists.txt b/example/01_gemm/CMakeLists.txt index db3498dca2..97ac21eba5 100755 --- a/example/01_gemm/CMakeLists.txt +++ b/example/01_gemm/CMakeLists.txt @@ -30,8 +30,6 @@ add_example_executable(example_gemm_xdl_fp8_v3 gemm_xdl_fp8_v3.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp8_v3) add_example_executable(example_gemm_xdl_fp16_fp8_v3 gemm_xdl_fp16_fp8_v3.cpp) add_example_executable(example_gemm_xdl_fp16_pk_i4_v3 gemm_xdl_fp16_pk_i4_v3.cpp) -add_example_executable(example_gemm_xdl_fp8_pk_i4_bpreshuffle_v3 gemm_xdl_fp8_pk_i4_bpreshuffle_v3.cpp) -add_example_executable(example_gemm_xdl_fp8_pk_i4_v3 gemm_xdl_fp8_pk_i4_v3.cpp) add_example_executable(example_gemm_xdl_fp16_pk_i4_v3_b_scale gemm_xdl_fp16_pk_i4_v3_b_scale.cpp) add_example_executable(example_gemm_xdl_bf16_pk_i4_v3 gemm_xdl_bf16_pk_i4_v3.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16_fp8_v3) diff --git a/example/01_gemm/common.hpp b/example/01_gemm/common.hpp index 58b15bda1d..9664c50b6e 100644 --- a/example/01_gemm/common.hpp +++ b/example/01_gemm/common.hpp @@ -369,26 +369,3 @@ inline __host__ __device__ constexpr double get_atol() return 1e-3; } } - -float i4_to_f32_gfx9(uint8_t i4) -{ - static std::unordered_map u = { - {0b1000, -0.5000f}, - {0b1001, -0.4375f}, - {0b1010, -0.3750f}, - {0b1011, -0.3125f}, - {0b1100, -0.2500f}, - {0b1101, -0.1875f}, - {0b1110, -0.1250f}, - {0b1111, -0.0625f}, - {0b0 , +0.0000f}, - {0b1 , +0.0625f}, - {0b10 , +0.1250f}, - {0b11 , +0.1875f}, - {0b100 , +0.2500f}, - {0b101 , +0.3125f}, - {0b110 , +0.3750f}, - {0b111 , +0.4375f}}; - - return u[i4]; -} diff --git a/example/01_gemm/gemm_xdl_fp8_pk_i4_bpreshuffle_v3.cpp b/example/01_gemm/gemm_xdl_fp8_pk_i4_bpreshuffle_v3.cpp deleted file mode 100644 index 544438bccb..0000000000 --- a/example/01_gemm/gemm_xdl_fp8_pk_i4_bpreshuffle_v3.cpp +++ /dev/null @@ -1,369 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. - -#include "common.hpp" - -#include "ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3_b_preshuffle.hpp" - -using F8 = ck::f8_t; -using I4 = ck::pk_i4_t; -using F16 = ck::half_t; -using F32 = float; - -using ADataType = F8; -using BDataType = I4; -using AccDataType = F32; -using CShuffleDataType = F16; -using CDataType = F16; - -using ALayout = Row; -using BLayout = Col; -using CLayout = Row; - -using AElementOp = PassThrough; -using BElementOp = PassThrough; -using CElementOp = PassThrough; - -static constexpr auto GemmDefault = ck::tensor_operation::device::GemmSpecialization::Default; - -static constexpr bool PermuteA = false; -static constexpr bool PermuteB = false; - -// clang-format off -#if 0 -using DeviceGemmV2Instance = - ck::tensor_operation::device::DeviceGemm_Xdl_CShuffleV3_BPreshuffle< - ALayout, BLayout, CLayout, - ADataType, BDataType, CDataType, AccDataType, CShuffleDataType, - AElementOp, BElementOp, CElementOp, GemmDefault, - 256, - 128, 128, - 256, 16, 32, - 32, 32, - 4, 1, - S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, - 2, 16, 16, 0, - S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, - 2, 32, 32, 0, - 1, 1, S<1, 32, 1, 8>, 4, - ck::BlockGemmPipelineScheduler::Intrawave, ck::BlockGemmPipelineVersion::v1, F8, F8, PermuteA, PermuteB>; - -#else -using DeviceGemmV2Instance = - ck::tensor_operation::device::DeviceGemm_Xdl_CShuffleV3_BPreshuffle< - ALayout, BLayout, CLayout, - ADataType, BDataType, CDataType, AccDataType, CShuffleDataType, - AElementOp, BElementOp, CElementOp, GemmDefault, - 256, - 256, 256, - 128, 16, 32, - 32, 32, - 4, 4, - S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, - 2, 16, 16, 0, - S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, - 2, 32, 32, 0, - 1, 1, S<1, 32, 1, 8>, 8, - ck::BlockGemmPipelineScheduler::Intrawave, ck::BlockGemmPipelineVersion::v3, F8, F8, PermuteA, PermuteB>; - -#endif -// clang-format on - -template -bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) -{ - using namespace ck::literals; - - auto M = problem_size.M; - auto N = problem_size.N; - auto K = problem_size.K; - auto StrideA = problem_size.StrideA; - auto StrideB = problem_size.StrideB; - auto StrideC = problem_size.StrideC; - auto KBatch = problem_size.KBatch; - - 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) - { - return HostTensorDescriptor({row, col}, {stride, 1_uz}); - } - else - { - return HostTensorDescriptor({row, col}, {1_uz, stride}); - } - }; - - auto f_get_default_stride = - [](std::size_t row, std::size_t col, ck::index_t stride, auto layout) { - if(stride == -1) - { - // give a chance if stride is -1, return a default packed stride - if constexpr(std::is_same_v) - { - return static_cast(col); - } - else - { - return static_cast(row); - } - } - else - 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{}); - - 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 b_k_n_preshuffled(f_host_tensor_descriptor(K, N, StrideB, BLayout{})); - - switch(config.init_method) - { - case 0: - a_m_k.GenerateTensorValue(GeneratorTensor_1{1}); - b_k_n.GenerateTensorValue(GeneratorTensor_1{1}); - break; - case 1: - a_m_k.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - b_k_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - break; - case 2: - a_m_k.GenerateTensorValue(GeneratorTensor_1{1}); - b_k_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - break; - case 3: - a_m_k.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - b_k_n.GenerateTensorValue(GeneratorTensor_1{1}); - break; - default: - a_m_k.GenerateTensorValue(GeneratorTensor_3{0.0, 1.0}); - b_k_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - } - - 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{})); - - std::cout << "a_m_k: " << a_m_k.mDesc << std::endl; - std::cout << "b_k_n: " << b_k_n.mDesc << std::endl; - std::cout << "b_k_n_preshuffled:" << b_k_n_preshuffled.mDesc << std::endl; - std::cout << "c_m_n: " << c_m_n_host_result.mDesc << std::endl; - - DeviceMem a_m_k_device_buf(sizeof(ADataType) * a_m_k.mDesc.GetElementSpaceSize()); - DeviceMem b_k_n_device_buf(sizeof(BDataType) * b_k_n_preshuffled.mDesc.GetElementSpaceSize()); - DeviceMem c_m_n_device_buf(sizeof(CDataType) * c_m_n_device_result.mDesc.GetElementSpaceSize()); - - // do GEMM - auto gemm = DeviceGemmV2Instance{}; - - // weight pre-shuffle - int KPack = 32; // int4 -> 32, fp8 -> 16, fp16 -> 8 - int NLane = gemm.GetPreShuffleParameters(); - int KLane = 64 / NLane; - - int K0 = K / (KLane * KPack); - // K -> K0 KLane KPack - // N -> N0 NLane - // N, K -> N0 K0 KLane NLane KPack - int tempk; - for(int n = 0; n < N; ++n) - { - for(int k = 0; k < K; ++k) - { - int n0 = n / NLane; - int n1 = n % NLane; - - int k0 = k / (KLane * KPack); - tempk = k % (KLane * KPack); - int k1 = tempk / KPack; - int k2 = tempk % KPack; - - int outputIndex = n0 * KPack * NLane * KLane * K0 + k0 * KPack * NLane * KLane + - k1 * KPack * NLane + n1 * KPack + k2; - - b_k_n_preshuffled(outputIndex) = b_k_n(n * K + k); - } - } - - // vector pk_i4x4 permute - for(int i = 0; i < N; i++) - { - for(int j = 0; j < K; j += 8) - { - int input[8]; - - for(int k = 0; k < 4; k++) - { - int i4x2 = b_k_n_preshuffled(j + k * 2, i).data; - input[k * 2 + 0] = (i4x2 >> 4) & 0xf; - input[k * 2 + 1] = (i4x2 >> 0) & 0xf; - } - - // permute 01234567->20643175 - { - int hi = input[2]; - int lo = input[0]; - int i4x2 = (hi << 4) | lo; - - b_k_n_preshuffled(j + 0, i) = i4x2; - } - - { - int hi = input[6]; - int lo = input[4]; - int i4x2 = (hi << 4) | lo; - - b_k_n_preshuffled(j + 2, i) = i4x2; - } - - { - int hi = input[3]; - int lo = input[1]; - int i4x2 = (hi << 4) | lo; - - b_k_n_preshuffled(j + 4, i) = i4x2; - } - - { - int hi = input[7]; - int lo = input[5]; - int i4x2 = (hi << 4) | lo; - - b_k_n_preshuffled(j + 6, i) = i4x2; - } - } - } - - a_m_k_device_buf.ToDevice(a_m_k.mData.data()); - b_k_n_device_buf.ToDevice(b_k_n_preshuffled.mData.data()); - DeviceMem workspace; - - auto a_element_op = AElementOp{}; - auto b_element_op = BElementOp{}; - auto c_element_op = CElementOp{}; - - auto invoker = gemm.MakeInvoker(); - float ave_time = 0; - - auto argument = gemm.MakeArgument(static_cast(a_m_k_device_buf.GetDeviceBuffer()), - static_cast(b_k_n_device_buf.GetDeviceBuffer()), - static_cast(c_m_n_device_buf.GetDeviceBuffer()), - M, - N, - K, - StrideA, - StrideB, - StrideC, - KBatch, - a_element_op, - b_element_op, - c_element_op); - - if(!gemm.IsSupportedArgument(argument)) - { - std::cerr << gemm.GetTypeString() << " does not support this problem" << std::endl; - - return true; - } - - bool pass = true; - if(config.do_verification) - { - Tensor b_k_n_f32({K, N}); - - for(int n = 0; n < N; n++) - { - for(int k = 0; k < K; k++) - { - ck::pk_i4_t i4x2 = b_k_n(k, n).data; - uint8_t i4 = 0; - - if(k % 2 == 1) - i4 = (i4x2.data >> 0) & 0xf; - else - i4 = (i4x2.data >> 4) & 0xf; - - float v_b = i4_to_f32_gfx9(i4); - b_k_n_f32(k, n) = v_b; - } - } - - using ReferenceGemmInstance = ck::tensor_operation::host::ReferenceGemm; - - auto ref_gemm = ReferenceGemmInstance{}; - auto ref_invoker = ref_gemm.MakeInvoker(); - - auto ref_argument = ref_gemm.MakeArgument( - a_m_k, b_k_n_f32, c_m_n_host_result, PassThrough{}, PassThrough{}, PassThrough{}); - - ref_invoker.Run(ref_argument); - - ave_time = invoker.Run(argument, StreamConfig{nullptr, false, 0}); - c_m_n_device_buf.FromDevice(c_m_n_device_result.mData.data()); - - pass &= ck::utils::check_err(c_m_n_device_result, - c_m_n_host_result, - "Error: Incorrect results!", - get_rtol(), - get_atol()); - } - - if(config.time_kernel) - { - ave_time = - invoker.Run(argument, StreamConfig{nullptr, config.time_kernel, 0, 20, 50, true, 50}); - - std::size_t flop = 2_uz * M * N * K; - std::size_t num_btype = - sizeof(ADataType) * M * K + - sizeof(BDataType) * K * N / - (ck::is_same_v, ck::pk_i4_t> ? 2 : 1) + - sizeof(CDataType) * M * N; - - float tflops = static_cast(flop) / 1.E9 / ave_time; - - float gb_per_sec = num_btype / 1.E6 / ave_time; - - std::cout << "Perf: " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec - << " GB/s, " << gemm.GetTypeString() << std::endl; - } - -#if 0 - printf("B Matrix:\n"); - for(int n = 0; n < N; n++) - { - for(int k = 0; k < K; k++) - { - ck::pk_i4_t i4x2 = b_k_n(k, n).data; - int8_t i4 = 0; - if(k % 2 == 1) - i4 = (i4x2.data >> 0) & 0xf; - else - i4 = (i4x2.data >> 4) & 0xf; - - printf("%f (%d),", i4_to_f32_gfx9(i4), static_cast(i4x2.data)); - } - printf("\n"); - } -#endif - - return pass; -} - -bool run_gemm_splitk_example(int argc, char* argv[]) -{ - ProblemSizeSplitK problem_size; - ExecutionConfig config; - - return parse_cmd_args(argc, argv, problem_size, config) && run_gemm(problem_size, config); -} - -int main(int argc, char* argv[]) { return !run_gemm_splitk_example(argc, argv); } diff --git a/example/01_gemm/gemm_xdl_fp8_pk_i4_v3.cpp b/example/01_gemm/gemm_xdl_fp8_pk_i4_v3.cpp deleted file mode 100644 index 4dd92b1973..0000000000 --- a/example/01_gemm/gemm_xdl_fp8_pk_i4_v3.cpp +++ /dev/null @@ -1,348 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - -#include "common.hpp" - -#include "ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp" - -using F8 = ck::f8_t; -using I4 = ck::pk_i4_t; -using F16 = ck::half_t; -using F32 = float; - -using ADataType = F8; -using BDataType = I4; -using AccDataType = float; -using CShuffleDataType = F16; -using CDataType = F16; - -using ALayout = Row; -using BLayout = Col; -using CLayout = Row; - -using AElementOp = PassThrough; -using BElementOp = PassThrough; -using CElementOp = PassThrough; - -static constexpr auto GemmDefault = ck::tensor_operation::device::GemmSpecialization::Default; - -static constexpr bool PermuteA = false; -static constexpr bool PermuteB = true; -static constexpr ck::index_t KPerBlock = 128; - -// clang-format off -using DeviceGemmV2Instance = - ck::tensor_operation::device::DeviceGemm_Xdl_CShuffleV3< - ALayout, BLayout, CLayout, - ADataType, BDataType, CDataType, AccDataType, CShuffleDataType, - AElementOp, BElementOp, CElementOp, GemmDefault, - 128, - 16, 128, - KPerBlock, 16, 32, - 16, 16, - 1, 4, - S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, - 2, 16, 16, 0, - S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, - 2, 32, 32, 0, - 1, 1, S<1, 16, 1, 8>, 4, - ck::BlockGemmPipelineScheduler::Interwave, ck::BlockGemmPipelineVersion::v2, ADataType, ADataType, PermuteA, PermuteB>; - -// clang-format on - -template -bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) -{ - using namespace ck::literals; - - auto M = problem_size.M; - auto N = problem_size.N; - auto K = problem_size.K; - auto StrideA = problem_size.StrideA; - auto StrideB = problem_size.StrideB; - auto StrideC = problem_size.StrideC; - auto KBatch = problem_size.KBatch; - - 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) - { - return HostTensorDescriptor({row, col}, {stride, 1_uz}); - } - else - { - return HostTensorDescriptor({row, col}, {1_uz, stride}); - } - }; - - auto f_get_default_stride = - [](std::size_t row, std::size_t col, ck::index_t stride, auto layout) { - if(stride == -1) - { - // give a chance if stride is -1, return a default packed stride - if constexpr(std::is_same_v) - { - return static_cast(col); - } - else - { - return static_cast(row); - } - } - else - 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{}); - - 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 b_k_n_permute(f_host_tensor_descriptor(K, N, StrideB, BLayout{})); - - switch(config.init_method) - { - case 0: - a_m_k.GenerateTensorValue(GeneratorTensor_1{1}); - b_k_n.GenerateTensorValue(GeneratorTensor_1{1}); - break; - case 1: - a_m_k.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - b_k_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - break; - case 2: - a_m_k.GenerateTensorValue(GeneratorTensor_1{1}); - b_k_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - break; - case 3: - a_m_k.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - b_k_n.GenerateTensorValue(GeneratorTensor_1{1}); - break; - default: - a_m_k.GenerateTensorValue(GeneratorTensor_3{0.0, 1.0}); - b_k_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - } - - 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{})); - - 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; - - DeviceMem a_m_k_device_buf(sizeof(ADataType) * a_m_k.mDesc.GetElementSpaceSize()); - DeviceMem b_k_n_device_buf(sizeof(BDataType) * b_k_n_permute.mDesc.GetElementSpaceSize()); - DeviceMem c_m_n_device_buf(sizeof(CDataType) * c_m_n_device_result.mDesc.GetElementSpaceSize()); - - // weight permute - if constexpr(PermuteB) - { - int K1 = KPerBlock; - int K0 = K / KPerBlock; - - // int K0, N, K1 - for(int j = 0; j < K0; j++) - { - for(int i = 0; i < N; i++) - { - for(int jj = 0; jj < K1; jj++) - { - b_k_n_permute(j * N * K1 + i * K1 + jj) = b_k_n(i * K + (j * K1 + jj)); - } - } - } - } - else - { - for(int i = 0; i < N; i++) - { - for(int j = 0; j < K; j++) - { - b_k_n_permute(i * K + j) = b_k_n(i * K + j); - } - } - } - - // vector pk_i4x4 permute - for(int i = 0; i < N; i++) - { - for(int j = 0; j < K; j += 8) - { - int input[8]; - - for(int k = 0; k < 4; k++) - { - int i4x2 = b_k_n_permute(j + k * 2, i).data; - input[k * 2 + 0] = (i4x2 >> 4) & 0xf; - input[k * 2 + 1] = (i4x2 >> 0) & 0xf; - } - - // permute 01234567->20643175 - { - int hi = input[2]; - int lo = input[0]; - int i4x2 = (hi << 4) | lo; - - b_k_n_permute(j + 0, i) = i4x2; - } - - { - int hi = input[6]; - int lo = input[4]; - int i4x2 = (hi << 4) | lo; - - b_k_n_permute(j + 2, i) = i4x2; - } - - { - int hi = input[3]; - int lo = input[1]; - int i4x2 = (hi << 4) | lo; - - b_k_n_permute(j + 4, i) = i4x2; - } - - { - int hi = input[7]; - int lo = input[5]; - int i4x2 = (hi << 4) | lo; - - b_k_n_permute(j + 6, i) = i4x2; - } - } - } - - a_m_k_device_buf.ToDevice(a_m_k.mData.data()); - b_k_n_device_buf.ToDevice(b_k_n_permute.mData.data()); - DeviceMem workspace; - - auto a_element_op = AElementOp{}; - auto b_element_op = BElementOp{}; - auto c_element_op = CElementOp{}; - - // do GEMM - auto gemm = DeviceGemmV2Instance{}; - auto invoker = gemm.MakeInvoker(); - float ave_time = 0; - - auto argument = gemm.MakeArgument(static_cast(a_m_k_device_buf.GetDeviceBuffer()), - static_cast(b_k_n_device_buf.GetDeviceBuffer()), - static_cast(c_m_n_device_buf.GetDeviceBuffer()), - M, - N, - K, - StrideA, - StrideB, - StrideC, - KBatch, - a_element_op, - b_element_op, - c_element_op); - - if(!gemm.IsSupportedArgument(argument)) - { - std::cerr << gemm.GetTypeString() << " does not support this problem" << std::endl; - - return true; - } - - bool pass = true; - if(config.do_verification) - { - Tensor b_k_n_f32({K, N}); - - for(int n = 0; n < N; n++) - { - for(int k = 0; k < K; k++) - { - ck::pk_i4_t i4x2 = b_k_n(k, n).data; - uint8_t i4 = 0; - - if(k % 2 == 1) - i4 = (i4x2.data >> 0) & 0xf; - else - i4 = (i4x2.data >> 4) & 0xf; - - float v_b = i4_to_f32_gfx9(i4); - b_k_n_f32(k, n) = v_b; - } - } - - using ReferenceGemmInstance = ck::tensor_operation::host::ReferenceGemm; - - auto ref_gemm = ReferenceGemmInstance{}; - auto ref_invoker = ref_gemm.MakeInvoker(); - - auto ref_argument = ref_gemm.MakeArgument( - a_m_k, b_k_n_f32, c_m_n_host_result, PassThrough{}, PassThrough{}, PassThrough{}); - - ref_invoker.Run(ref_argument); - - ave_time = invoker.Run(argument, StreamConfig{nullptr, false, 0}); - c_m_n_device_buf.FromDevice(c_m_n_device_result.mData.data()); - - pass &= ck::utils::check_err(c_m_n_device_result, - c_m_n_host_result, - "Error: Incorrect results!", - get_rtol(), - get_atol()); - } - - if(config.time_kernel) - { - ave_time = - invoker.Run(argument, StreamConfig{nullptr, config.time_kernel, 0, 20, 50, true, 50}); - - std::size_t flop = 2_uz * M * N * K; - std::size_t num_btype = - sizeof(ADataType) * M * K + - sizeof(BDataType) * K * N / - (ck::is_same_v, ck::pk_i4_t> ? 2 : 1) + - sizeof(CDataType) * M * N; - - float tflops = static_cast(flop) / 1.E9 / ave_time; - - float gb_per_sec = num_btype / 1.E6 / ave_time; - - std::cout << "Perf: " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec - << " GB/s, " << gemm.GetTypeString() << std::endl; - } - -#if 0 - printf("B Matrix:\n"); - for(int n = 0; n < N; n++) - { - for(int k = 0; k < K; k++) - { - ck::pk_i4_t i4x2 = b_k_n(k, n).data; - int8_t i4 = 0; - if(k % 2 == 1) - i4 = (i4x2.data >> 0) & 0xf; - else - i4 = (i4x2.data >> 4) & 0xf; - - printf("%f (%d),", i4_to_f32_gfx9(i4), static_cast(i4x2.data)); - } - printf("\n"); - } -#endif - - return pass; -} - -bool run_gemm_splitk_example(int argc, char* argv[]) -{ - ProblemSizeSplitK problem_size; - ExecutionConfig config; - - return parse_cmd_args(argc, argv, problem_size, config) && run_gemm(problem_size, config); -} - -int main(int argc, char* argv[]) { return !run_gemm_splitk_example(argc, argv); } diff --git a/example/65_gemm_multiply_multiply/moe_pk_i4_gemm1.cpp b/example/65_gemm_multiply_multiply/moe_pk_i4_gemm1.cpp deleted file mode 100644 index c02cc5efca..0000000000 --- a/example/65_gemm_multiply_multiply/moe_pk_i4_gemm1.cpp +++ /dev/null @@ -1,523 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - -#include -#include -#include -#include - -#include "ck/ck.hpp" -#include "ck/tensor_operation/gpu/device/gemm_specialization.hpp" -#include "ck/tensor_operation/gpu/device/impl/device_moe_gemm.hpp" -// #include "ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle_v3.hpp" -#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" -#include "ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp" - -#include "ck/library/utility/device_memory.hpp" -#include "ck/library/utility/host_tensor.hpp" -#include "ck/library/utility/host_tensor_generator.hpp" -#include "ck/library/utility/literals.hpp" -#include "ck/library/reference_tensor_operation/cpu/reference_moe_gemm.hpp" -#include "ck/library/utility/check_err.hpp" - -#include "ck/utility/blkgemmpipe_scheduler.hpp" - -template -using S = ck::Sequence; - -using I4 = ck::pk_i4_t; -using F16 = ck::half_t; -// using BF16 = ck::bhalf_t; -using F8 = ck::f8_t; -using F32 = float; - -using Row = ck::tensor_layout::gemm::RowMajor; -using Col = ck::tensor_layout::gemm::ColumnMajor; - -using A0DataType = F8; -using B0DataType = I4; -using EDataType = F16; -using AccDataType = F32; -using CShuffleDataType = F32; -using D0DataType = F32; -using D1DataType = F32; -using DsDataType = ck::Tuple; - -using A0Layout = Row; -using B0Layout = Col; -using ELayout = Row; -using D0Layout = Row; -using D1Layout = Col; -using DsLayout = ck::Tuple; - -// for gate, a_scale, b_scale -struct MulABScale -{ - template - __host__ __device__ constexpr void - operator()(E& e, const C& c, const D0& d0, const D1& d1) const; - - template <> - __host__ __device__ constexpr void operator() - (EDataType& e, - const float& c, - const float& d0, - const float& d1) const - { -#if CK_USE_PK4_LAYOUT_SHUFFLE - e = ck::type_convert(c * d1 * d0 * 16); -#else - e = ck::type_convert(c * d1 * d0); -#endif - } -}; - -// for gate, a_scale, b_scale, fuse silu, -struct MulABScaleSilu -{ - template - __host__ __device__ constexpr void - operator()(E& e, const C& c, const D0& d0, const D1& d1) const; - - template <> - __host__ __device__ constexpr void operator() - (EDataType& e, - const float& c, - const float& d0, - const float& d1) const - { - // act - float x0 = 0; -#if CK_USE_PK4_LAYOUT_SHUFFLE - ck::tensor_operation::element_wise::Silu{}(x0, c * d1 * d0 * 16); -#else - ck::tensor_operation::element_wise::Silu{}(x0, c * d1 * d0); -#endif - e = ck::type_convert(x0); - } -}; - -using CDEElementOp = MulABScale; - -#if 1 -void preShuffleBuffer(const I4* src, I4* dst, int N, int K, int NXdl) -{ - int KPack = 32; - int NLane = NXdl; - int KLane = 64 / NLane; - - int K0 = K / (KLane * KPack); - // K -> K0 KLane KPack - // N -> N0 NLane - // N, K -> N0 K0 KLane NLane KPack - int tempk; - for(int n = 0; n < N; ++n) - { - for(int k = 0; k < K; ++k) - { - int n0 = n / NLane; - int n1 = n % NLane; - - int k0 = k / (KLane * KPack); - tempk = k % (KLane * KPack); - int k1 = tempk / KPack; - int k2 = tempk % KPack; - - int outputIndex = n0 * KPack * NLane * KLane * K0 + k0 * KPack * NLane * KLane + - k1 * KPack * NLane + n1 * KPack + k2; - - dst[outputIndex / 2] = src[(n * K + k) / 2]; - } - } -} -#endif - -using PassThrough = ck::tensor_operation::element_wise::PassThrough; - -using AElementOp = PassThrough; -using BElementOp = PassThrough; - -static constexpr auto GemmSpec = ck::tensor_operation::device::GemmSpecialization::Default; -#if 0 -static constexpr ck::index_t MPerBlock = 64; -static constexpr ck::index_t MXDLPerWave = 1; -static constexpr ck::index_t NXDLPerWave = 2; -static constexpr ck::index_t BLOCKSIZE = 256; -static constexpr ck::index_t NPerBlock = 128; -static constexpr ck::index_t MNPerXDL = 32; -static constexpr ck::index_t KPerBlock = 64 / sizeof(A0DataType); -static constexpr ck::index_t Nswizzle = false; -static constexpr ck::index_t AK1 = 16 / sizeof(A0DataType); -static constexpr ck::index_t BK1 = 32 / sizeof(B0DataType); -static constexpr ck::index_t EVec = 16 / sizeof(EDataType); -static constexpr ck::index_t D0Vec = 1; -static constexpr ck::index_t D1Vec = 1; - -// clang-format off -using DeviceOpInstance = ck::tensor_operation::device::DeviceMoeGemm< - Row, Col, DsLayout, ELayout, - A0DataType, B0DataType, DsDataType, EDataType, AccDataType, CShuffleDataType, - AElementOp, BElementOp, CDEElementOp, GemmSpec, - BLOCKSIZE, MPerBlock, NPerBlock, KPerBlock, - AK1, BK1, - MNPerXDL, MNPerXDL, - MXDLPerWave, NXDLPerWave, - S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, AK1, AK1, 0, - S<2, 128, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, BK1, BK1, 0, - MXDLPerWave, 1, S<1, 32, 1, 8>, S, - ck::BlockGemmPipelineScheduler::Intrawave, ck::BlockGemmPipelineVersion::v1, Nswizzle, true, A0DataType>; -// clang-format on -#else -static constexpr ck::index_t MPerBlock = 128; -static constexpr ck::index_t Nswizzle = false; -// clang-format off -using DeviceOpInstance = ck::tensor_operation::device::DeviceMoeGemm< - Row, Col, DsLayout, ELayout, - A0DataType, B0DataType, DsDataType, EDataType, AccDataType, CShuffleDataType, - AElementOp, BElementOp, CDEElementOp, GemmSpec, - 256, MPerBlock, 128, 128, - 16, 32, - 32, 32, - 4, 1, - S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, - S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 32, 32, 0, - 2, 1, S<1, 32, 1, 8>, S<4, 1, 1>, - ck::BlockGemmPipelineScheduler::Intrawave, ck::BlockGemmPipelineVersion::v1, Nswizzle, true, A0DataType>; -// clang-format on -#endif - -int main(int argc, char* argv[]) -{ - bool do_verification = true; - int init_method = 1; - bool time_kernel = true; - -// tokens = 1 -// topk = 1 -// experts = 8 -// per expert: - // GEMM shape - ck::index_t N = 14336 * 2; - ck::index_t K = 4096; - ck::index_t experts = 8; - ck::index_t sorted_tile_num = 16; - ck::index_t valid_tile_num = 13; - ck::index_t sorted_size = sorted_tile_num * MPerBlock; - ck::index_t valid_size = valid_tile_num * MPerBlock; - ck::index_t tokens = 64; - ck::index_t topk = 2; - - // ck::index_t tokens = batch * topk; - - if(argc == 1) - { - // use default case - } - else if(argc == 7) - { - do_verification = std::stoi(argv[1]); - init_method = std::stoi(argv[2]); - time_kernel = std::stoi(argv[3]); - N = std::stoi(argv[4]); - K = std::stoi(argv[5]); - tokens = std::stoi(argv[6]); - } - else - { - printf("arg1: verification (0=no, 1=yes)\n"); - printf("arg2: initialization (0=no init, 1=integer value, 2=decimal value)\n"); - printf("arg3: time kernel (0=no, 1=yes)\n"); - printf( - "arg4 to 5: N, K, tokens\n"); - exit(0); - } - - if (tokens * topk > valid_size) - { - printf("err config, tokens * topk > valid_size\n"); - exit(-1); - } - ck::index_t StrideA = K; - ck::index_t StrideB = K; - ck::index_t StrideE = N; - constexpr ck::index_t NumDTensor = DsDataType::Size(); - constexpr auto StrideDs = std::array{0, 0}; - - ck::index_t KBatch = 1; - - // const ck::index_t experts = 8; - Tensor expert_ids(HostTensorDescriptor({sorted_tile_num}, {1})); - Tensor sorted_token_ids(HostTensorDescriptor({sorted_size}, {1})); - Tensor max_token_id(HostTensorDescriptor({1 + sorted_tile_num})); - max_token_id.mData = {valid_size, 2, 2, 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 0, 0, 0}; - int eids[] = {0, 0, 1, 2, 3, 3, 4, 4, 5, 5, 6, 7, 7, 3, 3, 3}; - for (int i = 0; i < sorted_tile_num; i++) { - expert_ids.mData[i] = eids[i]; - } - int token_per_tile = tokens * topk / valid_tile_num; - int tokenid = 0; - // sorted_token_ids.mData[0] = 0; - for (int i = 0; i < sorted_size; i++) { - int tile_off = i % MPerBlock; - if(tile_off < token_per_tile) - { - sorted_token_ids.mData[i] = (tokenid % tokens) | ((tokenid / tokens) << 24); - tokenid++; - } - else - { - sorted_token_ids.mData[i] = tokens; - } - } - expert_ids.savetxt("expert_ids.txt", "int"); - sorted_token_ids.savetxt("sorted_token_ids.txt", "int"); - Tensor a0_t_k(HostTensorDescriptor({tokens, K}, {K, 1})); - Tensor b0_e_n_k(HostTensorDescriptor({experts, K, N}, {N*K, 1, K})); - Tensor b0_preshuffled(HostTensorDescriptor({experts, K, N}, {N*K, 1, K})); - Tensor d0_t_n(HostTensorDescriptor({tokens, N}, {StrideDs[0], 0})); - Tensor d1_e_n(HostTensorDescriptor({experts, N}, {1, StrideDs[1]})); - Tensor e_t_n_host_result(HostTensorDescriptor({tokens, topk, N}, {topk * N, N, 1})); - Tensor e_t_n_device_result(HostTensorDescriptor({tokens, topk, N}, {topk * N, N, 1})); - - std::cout << "a0_t_k: " << a0_t_k.mDesc << std::endl; - std::cout << "b0_e_n_k: " << b0_e_n_k.mDesc << std::endl; - std::cout << "d1_e_n: " << d1_e_n.mDesc << std::endl; - std::cout << "d0_t_n: " << d0_t_n.mDesc << std::endl; - std::cout << "e_t_n: " << e_t_n_host_result.mDesc << std::endl; - - switch(init_method) - { - case 0: break; - case 1: - a0_t_k.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - b0_e_n_k.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - d0_t_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - d1_e_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - break; - case 2: - a0_t_k.GenerateTensorValue(GeneratorTensor_1{}); - b0_e_n_k.GenerateTensorValue(GeneratorTensor_1{}); - d0_t_n.GenerateTensorValue(GeneratorTensor_1{}); - d1_e_n.GenerateTensorValue(GeneratorTensor_1{}); - break; - default: - a0_t_k.GenerateTensorValue(GeneratorTensor_3{0.0, 1.0}); - b0_e_n_k.GenerateTensorValue(GeneratorTensor_3{-0.5, 0.5}); - d0_t_n.GenerateTensorValue(GeneratorTensor_3{0.0, 1.0}); - d1_e_n.GenerateTensorValue(GeneratorTensor_3{0.0, 1.0}); - } - DeviceMem sorted_token_ids_dev(sizeof(ck::index_t) * sorted_token_ids.mDesc.GetElementSpaceSize()); - DeviceMem expert_ids_dev(sizeof(ck::index_t) * expert_ids.mDesc.GetElementSpaceSize()); - DeviceMem max_token_id_dev(sizeof(ck::index_t) * max_token_id.mDesc.GetElementSpaceSize()); - DeviceMem a0_device_buf(sizeof(A0DataType) * a0_t_k.mDesc.GetElementSpaceSize()); - DeviceMem b0_device_buf(sizeof(B0DataType) * b0_e_n_k.mDesc.GetElementSpaceSize()); - DeviceMem d0_device_buf(sizeof(D0DataType) * d0_t_n.mDesc.GetElementSpaceSize()); - DeviceMem d1_device_buf(sizeof(D1DataType) * d1_e_n.mDesc.GetElementSpaceSize()); - DeviceMem e_device_buf(sizeof(EDataType) * e_t_n_device_result.mDesc.GetElementSpaceSize()); - a0_t_k.savetxt("a.txt"); - d0_t_n.savetxt("d0_t_n.txt", "int"); - d1_e_n.savetxt("d1_e_n.txt", "int"); - sorted_token_ids_dev.ToDevice(sorted_token_ids.mData.data()); - expert_ids_dev.ToDevice(expert_ids.mData.data()); - max_token_id_dev.ToDevice(max_token_id.mData.data()); - a0_device_buf.ToDevice(a0_t_k.mData.data()); - d0_device_buf.ToDevice(d0_t_n.mData.data()); - d1_device_buf.ToDevice(d1_e_n.mData.data()); - - auto a_element_op = AElementOp{}; - auto b_element_op = BElementOp{}; - auto cde_element_op = CDEElementOp{}; - - // do GEMM - auto device_op = DeviceOpInstance{}; - -#if 1 - preShuffleBuffer(b0_e_n_k.mData.data(), b0_preshuffled.mData.data(), N * experts, K, device_op.GetPreShuffleParameters()); -#else - // weight pre-shuffle - int KPack = 32; // int4 -> 32, fp8 -> 16, fp16 -> 8 - int NLane = device_op.GetPreShuffleParameters(); - int KLane = 64 / NLane; - - int K0 = K / (KLane * KPack); - // K -> K0 KLane KPack - // N -> N0 NLane - // N, K -> N0 K0 KLane NLane KPack - int tempk; - for(int e = 0; e < experts; ++e) - { - for(int n = 0; n < N; ++n) - { - for(int k = 0; k < K; ++k) - { - int n0 = n / NLane; - int n1 = n % NLane; - - int k0 = k / (KLane * KPack); - tempk = k % (KLane * KPack); - int k1 = tempk / KPack; - int k2 = tempk % KPack; - - int outputIndex = n0 * KPack * NLane * KLane * K0 + k0 * KPack * NLane * KLane + - k1 * KPack * NLane + n1 * KPack + k2; - - b0_preshuffled(e, outputIndex % K, outputIndex / K) = b0_e_n_k(e, k, n); - } - } - } -#endif - -#if CK_USE_PK4_LAYOUT_SHUFFLE - // vector pk_i4x4 permute - for(int e = 0; e < experts; e++) - { - for(int i = 0; i < N; i++) - { - for(int j = 0; j < K; j += 8) - { - int input[8]; - - for(int k = 0; k < 4; k++) - { - int i4x2 = b0_preshuffled(e, j + k * 2, i).data; - input[k * 2 + 0] = (i4x2 >> 4) & 0xf; - input[k * 2 + 1] = (i4x2 >> 0) & 0xf; - } - - // permute 01234567->20643175 - { - int hi = input[2]; - int lo = input[0]; - int i4x2 = (hi << 4) | lo; - - b0_preshuffled(e, j + 0, i) = i4x2; - } - - { - int hi = input[6]; - int lo = input[4]; - int i4x2 = (hi << 4) | lo; - - b0_preshuffled(e, j + 2, i) = i4x2; - } - - { - int hi = input[3]; - int lo = input[1]; - int i4x2 = (hi << 4) | lo; - - b0_preshuffled(e, j + 4, i) = i4x2; - } - - { - int hi = input[7]; - int lo = input[5]; - int i4x2 = (hi << 4) | lo; - - b0_preshuffled(e, j + 6, i) = i4x2; - } - } - } - } -#endif - - b0_device_buf.ToDevice(b0_preshuffled.mData.data()); - - auto invoker = device_op.MakeInvoker(); - auto argument = - device_op.MakeArgument(sorted_token_ids_dev.GetDeviceBuffer(), - expert_ids_dev.GetDeviceBuffer(), - max_token_id_dev.GetDeviceBuffer(), - a0_device_buf.GetDeviceBuffer(), - b0_device_buf.GetDeviceBuffer(), - std::array{d0_device_buf.GetDeviceBuffer(), - d1_device_buf.GetDeviceBuffer()}, - e_device_buf.GetDeviceBuffer(), - tokens, - topk, - sorted_size, - N, - K, - StrideA, - StrideB, - StrideDs, - StrideE, - KBatch, - a_element_op, - b_element_op, - cde_element_op); - - if(!device_op.IsSupportedArgument(argument)) - { - throw std::runtime_error( - "wrong! device_gemm with the specified compilation parameters does " - "not support this GEMM problem"); - } - if (time_kernel) { - float ave_time = invoker.Run(argument, StreamConfig{nullptr, time_kernel}); - - std::size_t flop = std::size_t(2) * tokens * topk * N * K; - std::size_t num_btype = - sizeof(A0DataType) * valid_tile_num * K + sizeof(B0DataType) / 2 * K * N * experts + sizeof(EDataType) * valid_tile_num * N; - - float tflops = static_cast(flop) / 1.E9 / ave_time; - - float gb_per_sec = num_btype / 1.E6 / ave_time; - - std::cout << "Perf: " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec - << " GB/s" << device_op.GetTypeString() << std::endl; - } - - if(do_verification) - { - invoker.Run(argument, StreamConfig{nullptr, false, 0 ,0,1}); - - e_device_buf.FromDevice(e_t_n_device_result.mData.data()); - - Tensor c_t_k_n({tokens, topk, N}, {topk * N, N, 1}); - - using ReferenceGemmInstance = ck::tensor_operation::host::ReferenceMoeGemm; - auto ref_moe_gemm = ReferenceGemmInstance{}; - auto ref_invoker = ref_moe_gemm.MakeInvoker(); - - auto ref_argument = ref_moe_gemm.MakeArgument( - sorted_token_ids, expert_ids, max_token_id, MPerBlock, a0_t_k, b0_e_n_k, c_t_k_n, PassThrough{}, PassThrough{}, PassThrough{}); - - ref_invoker.Run(ref_argument); - for(int m = 0; m < valid_size; ++m) - { - - const int fuse_t = sorted_token_ids.mData[m]; - const int t = fuse_t & 0xffffff; - const int topk_id = (fuse_t & 0xff000000) >> 24; - // printf("m %d fuset %d %d %d\n",m, fuse_t, t, topk_id); - - if (t >= tokens) - { - continue; - } - const int e = expert_ids(m / MPerBlock); - for(int n = 0; n < N; ++n) - { - cde_element_op(e_t_n_host_result(t, topk_id, n), c_t_k_n(t, topk_id, n), d0_t_n(t, n), d1_e_n(e, n)); - // printf("m %d n %d topk %d token %d %f %f\n",m, n,topk_id, t, e_t_n_host_result(t, topk_id, n), c_t_k_n(t, topk_id, n)); - } - } - - e_device_buf.FromDevice(e_t_n_device_result.mData.data()); - e_t_n_device_result.savetxt("out.txt"); - e_t_n_host_result.savetxt("ref.txt"); - return ck::utils::check_err( - e_t_n_device_result, e_t_n_host_result, "Error: Incorrect results!", 1e-3, 5e-2) - ? 0 - : 1; - } - - return 0; -} diff --git a/example/65_gemm_multiply_multiply/moe_pk_i4_gemm2.cpp b/example/65_gemm_multiply_multiply/moe_pk_i4_gemm2.cpp deleted file mode 100644 index 0a3e8dca1c..0000000000 --- a/example/65_gemm_multiply_multiply/moe_pk_i4_gemm2.cpp +++ /dev/null @@ -1,491 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - -#include -#include -#include -#include - -#include "ck/ck.hpp" -#include "ck/tensor_operation/gpu/device/gemm_specialization.hpp" -#include "ck/tensor_operation/gpu/device/impl/device_moe_gemm.hpp" -#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" -#include "ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp" - -#include "ck/library/utility/device_memory.hpp" -#include "ck/library/utility/host_tensor.hpp" -#include "ck/library/utility/host_tensor_generator.hpp" -#include "ck/library/utility/literals.hpp" -#include "ck/library/reference_tensor_operation/cpu/reference_moe_gemm2.hpp" -#include "ck/library/utility/check_err.hpp" - -#include "ck/utility/blkgemmpipe_scheduler.hpp" - -template -using S = ck::Sequence; - -using I4 = ck::pk_i4_t; -using F16 = ck::half_t; -using F8 = ck::f8_t; -using F32 = float; - -using Row = ck::tensor_layout::gemm::RowMajor; -using Col = ck::tensor_layout::gemm::ColumnMajor; - -using A0DataType = F8; -using B0DataType = I4; -using EDataType = F16; -using AccDataType = F32; -using CShuffleDataType = F32; -using D0DataType = F32; -using D1DataType = F32; -using D2DataType = F32; -using DsDataType = ck::Tuple; - -using A0Layout = Row; -using B0Layout = Col; -using ELayout = Row; -using D0Layout = Row; -using D1Layout = Col; -using D2Layout = ELayout; -// using DsLayoutGate = ck::Tuple; -using DsLayout = ck::Tuple; - -// d0: ascale, d1: bscale, d2:expert weight -struct MulABScaleExpertWeight -{ - template - __host__ __device__ constexpr void - operator()(E& e, const C& c, const D0& d0, const D1& d1, const D2& d2) const; - //for real kernel use - template <> - __host__ __device__ constexpr void operator() - (EDataType& e, - const float& c, - const float& d0, - const float& d1, - const float& d2) const - { - //for real kernel use - //warning: hack hack hack here!!!! ignore d0 right now as kernel mul d0 * d2 outside. tofix:felix - (void) d0; - -#if CK_USE_PK4_LAYOUT_SHUFFLE - e = ck::type_convert(c * d1 * d2 * 16); -#else - e = ck::type_convert(c * d1 * d2); -#endif - } - // for reference cpu - template <> - __host__ __device__ constexpr void operator() - (float& e, - const float& c, - const float& d0, - const float& d1, - const float& d2) const - { - // for reference cpu -#if CK_USE_PK4_LAYOUT_SHUFFLE - e = ck::type_convert(c * d0 * d1 * d2 * 16); -#else - e = ck::type_convert(c * d0 * d1 * d2); -#endif - } -}; - -using CDEElementOp = MulABScaleExpertWeight; - -void preShuffleBuffer(const I4* src, I4* dst, int N, int K, int NXdl) -{ - int KPack = 32; - int NLane = NXdl; - int KLane = 64 / NLane; - - int K0 = K / (KLane * KPack); - // K -> K0 KLane KPack - // N -> N0 NLane - // N, K -> N0 K0 KLane NLane KPack - int tempk; - for(int n = 0; n < N; ++n) - { - for(int k = 0; k < K; ++k) - { - int n0 = n / NLane; - int n1 = n % NLane; - - int k0 = k / (KLane * KPack); - tempk = k % (KLane * KPack); - int k1 = tempk / KPack; - int k2 = tempk % KPack; - - int outputIndex = n0 * KPack * NLane * KLane * K0 + k0 * KPack * NLane * KLane + - k1 * KPack * NLane + n1 * KPack + k2; - - dst[outputIndex / 2] = src[(n * K + k) / 2]; - } - } -} - -using PassThrough = ck::tensor_operation::element_wise::PassThrough; - -using AElementOp = PassThrough; -using BElementOp = PassThrough; -using CDEElementOp = MulABScaleExpertWeight; - -static constexpr auto GemmSpec = ck::tensor_operation::device::GemmSpecialization::Default; -static constexpr ck::index_t MPerBlock = 128; -static constexpr ck::index_t BLOCKSIZE = 256; -static constexpr ck::index_t MXDLPerWave = 2; -static constexpr ck::index_t NXDLPerWave = 2; -static constexpr ck::index_t NPerBlock = 128; -static constexpr ck::index_t MNPerXDL = 32; -static constexpr ck::index_t KPerBlock = 128 / sizeof(A0DataType); -// static constexpr ck::index_t MXDLPerWave = MPerBlock / 32; //todo fix this constraint -// static constexpr ck::index_t CShuffleMXDLPerWave = MPerBlock / 32; -static constexpr ck::index_t CShuffleNLane = 32; -static constexpr ck::index_t CShuffleMLane = BLOCKSIZE / CShuffleNLane; -static constexpr ck::index_t AK1 = 16 / sizeof(A0DataType); -static constexpr ck::index_t BK1 = 32 / sizeof(B0DataType); -static constexpr ck::index_t EVec = 2; -static constexpr ck::index_t D0Vec = 1; -static constexpr ck::index_t D1Vec = 1; -static constexpr ck::index_t D2Vec = 1; -using DeviceOpInstance = ck::tensor_operation::device::DeviceMoeGemm - // clang-format off - < Row, Col, DsLayout, ELayout, A0DataType, B0DataType, DsDataType, EDataType, AccDataType, CShuffleDataType, - AElementOp, BElementOp, CDEElementOp, GemmSpec, - BLOCKSIZE, MPerBlock, NPerBlock, KPerBlock, - AK1, BK1, - MNPerXDL, MNPerXDL, - MXDLPerWave, NXDLPerWave, - S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, AK1, AK1, 0, - S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, BK1, BK1, 0, - MXDLPerWave, 1, S<1, CShuffleMLane, 1, CShuffleNLane>, S, - ck::BlockGemmPipelineScheduler::Intrawave, ck::BlockGemmPipelineVersion::v1, false, false, A0DataType>; -// clang-format on - -int main(int argc, char* argv[]) -{ - bool do_verification = true; - int init_method = 1; - bool time_kernel = true; - -// tokens = 1 -// topk = 1 -// experts = 8 -// per expert: - // GEMM shape - ck::index_t N = 4096; - ck::index_t K = 14336; - ck::index_t experts = 8; - ck::index_t sorted_tile_num = 16; - ck::index_t valid_tile_num = 13; - ck::index_t sorted_size = sorted_tile_num * MPerBlock; - ck::index_t valid_size = valid_tile_num * MPerBlock; - ck::index_t tokens = 512; - ck::index_t topk = 2; - - if(argc == 1) - { - // use default case - } - else if(argc == 3) - { - // use default case - do_verification = std::stoi(argv[1]); - init_method = std::stoi(argv[2]); - time_kernel = std::stoi(argv[3]); - } - else if(argc == 7) - { - do_verification = std::stoi(argv[1]); - init_method = std::stoi(argv[2]); - time_kernel = std::stoi(argv[3]); - N = std::stoi(argv[4]); - K = std::stoi(argv[5]); - tokens = std::stoi(argv[6]); - } - else - { - printf("arg1: verification (0=no, 1=yes)\n"); - printf("arg2: initialization (0=no init, 1=integer value, 2=decimal value)\n"); - printf("arg3: time kernel (0=no, 1=yes)\n"); - printf( - "arg4 to 6: N, K, tokens\n"); - exit(0); - } - - ck::index_t StrideA = K; - ck::index_t StrideB = K; - ck::index_t StrideE = N; - constexpr ck::index_t NumDTensor = DsDataType::Size(); - constexpr auto StrideDs = std::array{0, 0, 0}; - - ck::index_t KBatch = 1; - - // const ck::index_t experts = 8; - Tensor expert_ids(HostTensorDescriptor({sorted_tile_num}, {1})); - Tensor sorted_token_ids(HostTensorDescriptor({sorted_size}, {1})); - Tensor max_token_id(HostTensorDescriptor({1})); - max_token_id.mData[0] = valid_size; - int eids[] = {0, 0, 1, 2, 3, 3, 4, 4, 5, 5, 6, 7, 7, 3, 3, 3}; - for (int i = 0; i < sorted_tile_num; i++) { - expert_ids.mData[i] = eids[i]; - } - if (tokens * topk > valid_size) - { - printf("err config, tokens * topk > valid_size\n"); - exit(-1); - } - int token_per_tile = tokens * topk / valid_tile_num; - int tokenid = 0; - // sorted_token_ids.mData[0] = 0; - for (int i = 0; i < sorted_size; i++) { - int tile_off = i % MPerBlock; - if(tile_off < token_per_tile) - { - sorted_token_ids.mData[i] = (tokenid % tokens) | ((tokenid / tokens) << 24); - tokenid++; - } - else - { - sorted_token_ids.mData[i] = tokens; - } - - } - expert_ids.savetxt("expert_ids.txt", "int"); - sorted_token_ids.savetxt("sorted_token_ids.txt", "int"); - Tensor a0_t_k_k(HostTensorDescriptor({tokens, topk, K}, {topk*K, K, 1})); - Tensor b0_e_n_k(HostTensorDescriptor({experts, K, N}, {N*K, 1, K})); - Tensor b0_preshuffled(HostTensorDescriptor({experts, K, N}, {N*K, 1, K})); - Tensor d0_t_n(HostTensorDescriptor({tokens, N}, {StrideDs[0], 0})); - Tensor d1_e_n(HostTensorDescriptor({experts, N}, {1, StrideDs[1]})); - Tensor d2_e_n(HostTensorDescriptor({sorted_size, N}, {1, 0})); - Tensor e_t_n_host_result(HostTensorDescriptor({tokens, N}, {N, 1})); - Tensor e_t_n_device_result(HostTensorDescriptor({tokens, N}, {N, 1})); - e_t_n_device_result.SetZero(); - std::cout << "a0_t_k_k: " << a0_t_k_k.mDesc << std::endl; - std::cout << "b0_e_n_k: " << b0_e_n_k.mDesc << std::endl; - std::cout << "d2_e_n: " << d2_e_n.mDesc << std::endl; - std::cout << "d1_e_n: " << d1_e_n.mDesc << std::endl; - std::cout << "d0_t_n: " << d0_t_n.mDesc << std::endl; - std::cout << "e_t_n: " << e_t_n_host_result.mDesc << std::endl; - - switch(init_method) - { - case 0: break; - case 1: - a0_t_k_k.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - b0_e_n_k.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - d0_t_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - d1_e_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - d2_e_n.GenerateTensorValue(GeneratorTensor_2{-2, 2}); - break; - case 2: - a0_t_k_k.GenerateTensorValue(GeneratorTensor_1{}); - b0_e_n_k.GenerateTensorValue(GeneratorTensor_1{}); - d0_t_n.GenerateTensorValue(GeneratorTensor_1{}); - d1_e_n.GenerateTensorValue(GeneratorTensor_1{}); - d2_e_n.GenerateTensorValue(GeneratorTensor_1{}); - break; - default: - a0_t_k_k.GenerateTensorValue(GeneratorTensor_3{0.0, 1.0}); - b0_e_n_k.GenerateTensorValue(GeneratorTensor_3{-0.5, 0.5}); - d0_t_n.GenerateTensorValue(GeneratorTensor_3{0.0, 1.0}); - d1_e_n.GenerateTensorValue(GeneratorTensor_3{0.0, 1.0}); - d2_e_n.GenerateTensorValue(GeneratorTensor_3{0.0, 1.0}); - } - DeviceMem sorted_token_ids_dev(sizeof(ck::index_t) * sorted_token_ids.mDesc.GetElementSpaceSize()); - DeviceMem expert_ids_dev(sizeof(ck::index_t) * expert_ids.mDesc.GetElementSpaceSize()); - DeviceMem max_token_id_dev(sizeof(ck::index_t) * max_token_id.mDesc.GetElementSpaceSize()); - DeviceMem a0_device_buf(sizeof(A0DataType) * a0_t_k_k.mDesc.GetElementSpaceSize()); - DeviceMem b0_device_buf(sizeof(B0DataType) * b0_e_n_k.mDesc.GetElementSpaceSize()); - DeviceMem d0_device_buf(sizeof(D0DataType) * d0_t_n.mDesc.GetElementSpaceSize()); - DeviceMem d1_device_buf(sizeof(D1DataType) * d1_e_n.mDesc.GetElementSpaceSize()); - DeviceMem d2_device_buf(sizeof(D2DataType) * d2_e_n.mDesc.GetElementSpaceSize()); - DeviceMem e_device_buf(sizeof(EDataType) * e_t_n_device_result.mDesc.GetElementSpaceSize()); - a0_t_k_k.savetxt("a.txt"); - expert_ids.savetxt("expert_ids.txt", "int"); - sorted_token_ids.savetxt("sorted_token_ids.txt", "int"); - d0_t_n.savetxt("d0_t_n.txt", "int"); - d1_e_n.savetxt("d1_e_n.txt", "int"); - d2_e_n.savetxt("d2_e_n.txt", "int"); - sorted_token_ids_dev.ToDevice(sorted_token_ids.mData.data()); - expert_ids_dev.ToDevice(expert_ids.mData.data()); - max_token_id_dev.ToDevice(max_token_id.mData.data()); - a0_device_buf.ToDevice(a0_t_k_k.mData.data()); - d0_device_buf.ToDevice(d0_t_n.mData.data()); - d1_device_buf.ToDevice(d1_e_n.mData.data()); - d2_device_buf.ToDevice(d2_e_n.mData.data()); - e_device_buf.ToDevice(e_t_n_device_result.mData.data()); - - auto a_element_op = AElementOp{}; - auto b_element_op = BElementOp{}; - auto cde_element_op = CDEElementOp{}; - - // do GEMM - auto device_op = DeviceOpInstance{}; - - preShuffleBuffer(b0_e_n_k.mData.data(), b0_preshuffled.mData.data(), N * experts, K, device_op.GetPreShuffleParameters()); - -#if CK_USE_PK4_LAYOUT_SHUFFLE - // vector pk_i4x4 permute - for(int e = 0; e < experts; e++) - { - for(int i = 0; i < N; i++) - { - for(int j = 0; j < K; j += 8) - { - int input[8]; - - for(int k = 0; k < 4; k++) - { - int i4x2 = b0_preshuffled(e, j + k * 2, i).data; - input[k * 2 + 0] = (i4x2 >> 4) & 0xf; - input[k * 2 + 1] = (i4x2 >> 0) & 0xf; - } - - // permute 01234567->20643175 - { - int hi = input[2]; - int lo = input[0]; - int i4x2 = (hi << 4) | lo; - - b0_preshuffled(e, j + 0, i) = i4x2; - } - - { - int hi = input[6]; - int lo = input[4]; - int i4x2 = (hi << 4) | lo; - - b0_preshuffled(e, j + 2, i) = i4x2; - } - - { - int hi = input[3]; - int lo = input[1]; - int i4x2 = (hi << 4) | lo; - - b0_preshuffled(e, j + 4, i) = i4x2; - } - - { - int hi = input[7]; - int lo = input[5]; - int i4x2 = (hi << 4) | lo; - - b0_preshuffled(e, j + 6, i) = i4x2; - } - } - } - } -#endif - - b0_device_buf.ToDevice(b0_preshuffled.mData.data()); - - auto invoker = device_op.MakeInvoker(); - auto argument = - device_op.MakeArgument(sorted_token_ids_dev.GetDeviceBuffer(), - expert_ids_dev.GetDeviceBuffer(), - max_token_id_dev.GetDeviceBuffer(), - a0_device_buf.GetDeviceBuffer(), - b0_device_buf.GetDeviceBuffer(), - std::array{d0_device_buf.GetDeviceBuffer(), - d1_device_buf.GetDeviceBuffer(), - d2_device_buf.GetDeviceBuffer()}, - e_device_buf.GetDeviceBuffer(), - tokens, - topk, - sorted_size, - N, - K, - StrideA, - StrideB, - StrideDs, - StrideE, - KBatch, - a_element_op, - b_element_op, - cde_element_op); - - if(!device_op.IsSupportedArgument(argument)) - { - throw std::runtime_error( - "wrong! device_gemm with the specified compilation parameters does " - "not support this GEMM problem"); - } - if (time_kernel) { - // not result correct here because output buf not setzero - float ave_time = invoker.Run(argument, StreamConfig{nullptr, time_kernel}); - - std::size_t flop = std::size_t(2) * tokens * topk * N * K; - std::size_t num_btype = - sizeof(A0DataType) * tokens * K * topk + sizeof(B0DataType) / 2 * K * N * experts + sizeof(EDataType) * tokens * N; - - float tflops = static_cast(flop) / 1.E9 / ave_time; - - float gb_per_sec = num_btype / 1.E6 / ave_time; - - std::cout << "Perf: " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec - << " GB/s" << device_op.GetTypeString() << std::endl; - } - - if(do_verification) - { - //gemm2 use atomic, so need to reinit outputs - e_device_buf.ToDevice(e_t_n_device_result.mData.data()); - invoker.Run(argument, StreamConfig{nullptr, false, 0 ,0,1}); - - Tensor c_t_n({tokens, N}); - - using ReferenceGemmInstance = - ck::tensor_operation::host::ReferenceMoeGemm2; - - auto ref_moe_gemm = ReferenceGemmInstance{}; - auto ref_invoker = ref_moe_gemm.MakeInvoker(); - auto ref_argument = ref_moe_gemm.MakeArgument(sorted_token_ids, - expert_ids, - max_token_id, - MPerBlock, - a0_t_k_k, - b0_e_n_k, - d0_t_n, - d1_e_n, - d2_e_n, - c_t_n, - PassThrough{}, - PassThrough{}, - cde_element_op); - - ref_invoker.Run(ref_argument); - for(int t = 0; t < tokens; ++t) - { - for(int n = 0; n < N; ++n) - { - e_t_n_host_result(t, n) = ck::type_convert(c_t_n(t, n)); - } - } - - e_device_buf.FromDevice(e_t_n_device_result.mData.data()); - e_t_n_device_result.savetxt("out.txt"); - e_t_n_host_result.savetxt("ref.txt"); - - return ck::utils::check_err( - e_t_n_device_result, e_t_n_host_result, "Error: Incorrect results!", 1e-3, 5e-2) - ? 0 - : 1; - } - - return 0; -} diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v1.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v1.hpp deleted file mode 100644 index d7b2de7d01..0000000000 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v1.hpp +++ /dev/null @@ -1,547 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. - -#pragma once - -#include "ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_base.hpp" - -namespace ck { - -// Compute optimized pipeline -// GlobalPrefetchStages: 2 -// LocalPreFillStages: 1 -// LocalPreFetchStages: 1 -// LocalSharedMemoryBuffer: 1 - -template -struct BlockwiseGemmXdlops_pipeline_bpreshuffle_bdequant_v1 -{ -}; - -template -struct BlockwiseGemmXdlops_pipeline_bpreshuffle_bdequant_v1 - : BlockwiseGemmXdlops_pipeline_base - -{ - using Base = BlockwiseGemmXdlops_pipeline_base; - using Base::A_K1; - using Base::B_K1; - using Base::I0; - using Base::I1; - using Base::KRepeat; - using Base::xdlops_gemm; - using typename Base::HotLoopInstList; - - using Base::a_block_desc_m0_m1_m2_k; - using Base::CalculateCThreadOriginDataIndex; - using Base::CalculateCThreadOriginDataIndex8D; - using Base::GetCBlockDescriptor_G_M0_N0_M1_N1_M2_M3_M4_N2; - using Base::GetCBlockDescriptor_M0_N0_M1_N1_M2_M3_M4_N2; - using Base::GetCBlockDescriptor_M0_N0_M1_N1_M2_N2_N3_N4; - using Base::GetCThreadBuffer; - using Base::GetCThreadDescriptor_G_M0_N0_M1_N1_M2_M3_M4_N2; - using Base::GetCThreadDescriptor_M0_N0_M1_N1_M2_M3_M4_N2; - using Base::GetCThreadDescriptor_M0_N0_M1_N1_M2_N2_N3_N4; - using Base::MakeCGridDescriptor_G_M0_N0_M1_N1_M2_M3_M4_N2; - using Base::MakeCGridDescriptor_M0_N0_M1_N1_M2_M3_M4_N2; - - using Base::AMmaKStride; - using Base::BMmaKStride; - - static constexpr index_t PrefetchStages = 2; - static constexpr index_t PrefillStages = 1; - static constexpr index_t GlobalBufferNum = 2; - - template - __host__ __device__ static constexpr auto MakeAGemmMmaTileDescriptor(const TileDesc_M0_M1_M2_K&) - { - constexpr index_t M0 = TileDesc_M0_M1_M2_K{}.GetLength(Number<0>{}); - constexpr index_t M1 = TileDesc_M0_M1_M2_K{}.GetLength(Number<1>{}); - constexpr index_t M2 = TileDesc_M0_M1_M2_K{}.GetLength(Number<2>{}); - constexpr index_t K2 = KPack; - constexpr index_t K1 = 64 / NPerXDL; - constexpr index_t K0 = KRepeat; - - return transform_tensor_descriptor( - TileDesc_M0_M1_M2_K{}, - make_tuple( - make_pass_through_transform(Number{}), - make_pass_through_transform(Number{}), - make_pass_through_transform(Number{}), - make_unmerge_transform(make_tuple(Number{}, Number{}, Number{}))), - make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}), - make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3, 4, 5>{})); - } - - static constexpr auto a_block_desc_m0_m1_m2_k0_k1_k2 = - MakeAGemmMmaTileDescriptor(a_block_desc_m0_m1_m2_k); - - __host__ __device__ static constexpr bool BlockHasHotloop(index_t num_loop) - { - return num_loop > PrefetchStages; - } - - __host__ __device__ static constexpr TailNumber BlockLoopTailNum(index_t num_loop) - { - return num_loop % 2 == 0 ? TailNumber::Even : TailNumber::Odd; - } - - __device__ static constexpr auto HotLoopScheduler() - { - constexpr auto num_ds_read_inst_a = HotLoopInstList::A_LDS_Read_Inst_Num; - constexpr auto num_buffer_load_inst_a = HotLoopInstList::A_Buffer_Load_Inst_Num; - constexpr auto num_buffer_load_inst_b = HotLoopInstList::B_Buffer_Load_Inst_Num; - - // B global - static_for<0, num_buffer_load_inst_b, 1>{}([&](auto i) { - ignore = i; - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read - }); - - // A global - static_for<0, num_buffer_load_inst_a, 1>{}([&](auto i) { - ignore = i; - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS write - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read - }); - - // A local - static_for<0, num_ds_read_inst_a / 2, 1>{}([&](auto i) { - ignore = i; - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 2, 0); // DS read - }); - } - - template - __device__ void Run(const AGridDesc& a_grid_desc, - const ABlockDesc& a_block_desc, - ABlockTransfer& a_blockwise_copy, - const AGridBuffer& a_grid_buf, - ABlockBuffer& a_block_buf, - const ABlockTransferStep& a_block_copy_step, - const BGridDesc& b_grid_desc, - BBlockTransfer& b_blockwise_copy, - const BGridBuffer& b_grid_buf, - BBlockBuffer& b_block_buf, - const BBlockTransferStep& b_block_copy_step, - CThreadBuffer& c_thread_buf, - index_t num_loop) const - { - ignore = b_block_buf; - __builtin_amdgcn_sched_barrier(0); - auto a_thread_buf = make_static_buffer( - a_thread_desc_.GetElementSpaceSize()); - auto b_thread_buf = make_static_buffer( - b_thread_desc_.GetElementSpaceSize()); - - auto b_thread_dequant_buf = make_static_buffer( - b_thread_desc_.GetElementSpaceSize()); - - StaticallyIndexedArray{}> b_thread_bufs; - constexpr auto b_block_origin_idx = make_tuple(I0, I0, I0, I0); - - StaticallyIndexedArray{}> b_thread_dequant_bufs; - - // Global prefetch A1 B1 - a_blockwise_copy.RunRead(a_grid_desc, a_grid_buf, I0); - b_blockwise_copy.Run(b_grid_desc, - b_grid_buf, - b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(I0)); - - a_blockwise_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step); - b_blockwise_copy.MoveSrcSliceWindow(b_grid_desc, b_block_copy_step); - __builtin_amdgcn_sched_barrier(0); - - // // Local prefill A1 - a_blockwise_copy.RunWrite(a_block_desc, a_block_buf, I0); - - // // Global prefetch A2 - a_blockwise_copy.RunRead(a_grid_desc, a_grid_buf, I0); - a_blockwise_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step); - - // Local prefetch A1 - block_sync_lds(); - static_for<0, MRepeat, 1>{}([&](auto m0) { - static_for<0, KRepeat, 1>{}([&](auto k0) { - a_thread_copy_.Run(a_block_desc_m0_m1_m2_k0_k1_k2, - make_tuple(m0, I0, I0, k0, I0, I0), - a_block_buf, - a_thread_desc_, - make_tuple(m0, I0, I0, k0, I0, I0), - a_thread_buf); - }); - }); - // B VGPR->VGPR dequant - b_thread_dequant_copy_.Run(b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(I0), - b_thread_desc_, - make_tuple(I0, I0, I0, I0), - b_thread_dequant_bufs(I0)); - - // Initialize C - c_thread_buf.Clear(); - - __builtin_amdgcn_sched_barrier(0); - - // main body - if constexpr(HasMainLoop) - { - index_t i = 0; - do - { - auto LoopFunc = [&](auto mfma_reg_buf, auto local_read_buf) { - b_blockwise_copy.Run(b_grid_desc, - b_grid_buf, - b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(local_read_buf)); - b_blockwise_copy.MoveSrcSliceWindow(b_grid_desc, b_block_copy_step); - - block_sync_lds(); - a_blockwise_copy.RunWrite(a_block_desc, a_block_buf, mfma_reg_buf); - - a_blockwise_copy.RunRead(a_grid_desc, a_grid_buf, local_read_buf); - a_blockwise_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step); - - static_for<0, MRepeat, 1>{}([&](auto m0) { - static_for<0, NRepeat, 1>{}([&](auto n0) { - static_for<0, KRepeat, 1>{}([&](auto k0) { - vector_type a_thread_vec; - vector_type b_thread_vec; - - static_for<0, KPack, 1>{}([&](auto ik) { - a_thread_vec.template AsType()(ik) = - a_thread_buf[Number{}]; - b_thread_vec.template AsType()(ik) = - b_thread_dequant_bufs[mfma_reg_buf] - [Number{}]; - }); - using mfma_input_type = - typename vector_type::type; - - constexpr index_t c_offset = - c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0)); - - xdlops_gemm.Run( - a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf.GetVectorTypeReference(Number{})); - }); - }); - }); - - block_sync_lds(); - - static_for<0, MRepeat, 1>{}([&](auto m0) { - static_for<0, KRepeat, 1>{}([&](auto k0) { - a_thread_copy_.Run(a_block_desc_m0_m1_m2_k0_k1_k2, - make_tuple(m0, I0, I0, k0, I0, I0), - a_block_buf, - a_thread_desc_, - make_tuple(m0, I0, I0, k0, I0, I0), - a_thread_buf); - }); - }); - // B VGPR->VGPR dequant - b_thread_dequant_copy_.Run(b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(local_read_buf), - b_thread_desc_, - make_tuple(I0, I0, I0, I0), - b_thread_dequant_bufs(local_read_buf)); - - HotLoopScheduler(); - __builtin_amdgcn_sched_barrier(0); - }; - - LoopFunc(I0, I1); - LoopFunc(I1, I0); - - i += 2; - } while(i < (num_loop - 2)); - } - // tail - if constexpr(TailNum == TailNumber::Even) - { - b_blockwise_copy.Run(b_grid_desc, - b_grid_buf, - b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(I1)); - - block_sync_lds(); - a_blockwise_copy.RunWrite(a_block_desc, a_block_buf); - - static_for<0, MRepeat, 1>{}([&](auto m0) { - static_for<0, NRepeat, 1>{}([&](auto n0) { - static_for<0, KRepeat, 1>{}([&](auto k0) { - vector_type a_thread_vec; - vector_type b_thread_vec; - - static_for<0, KPack, 1>{}([&](auto ik) { - a_thread_vec.template AsType()(ik) = - a_thread_buf[Number{}]; - b_thread_vec.template AsType()(ik) = - b_thread_dequant_bufs[I0][Number{}]; - }); - - using mfma_input_type = - typename vector_type::type; - - constexpr index_t c_offset = - c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0)); - - xdlops_gemm.Run(a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf.GetVectorTypeReference(Number{})); - }); - }); - }); - - block_sync_lds(); - - static_for<0, MRepeat, 1>{}([&](auto m0) { - static_for<0, KRepeat, 1>{}([&](auto k0) { - a_thread_copy_.Run(a_block_desc_m0_m1_m2_k0_k1_k2, - make_tuple(m0, I0, I0, k0, I0, I0), - a_block_buf, - a_thread_desc_, - make_tuple(m0, I0, I0, k0, I0, I0), - a_thread_buf); - }); - }); - // B VGPR->VGPR dequant - b_thread_dequant_copy_.Run(b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(I1), - b_thread_desc_, - make_tuple(I0, I0, I0, I0), - b_thread_dequant_bufs(I1)); - - __builtin_amdgcn_sched_barrier(0); - - static_for<0, MRepeat, 1>{}([&](auto m0) { - static_for<0, NRepeat, 1>{}([&](auto n0) { - static_for<0, KRepeat, 1>{}([&](auto k0) { - vector_type a_thread_vec; - vector_type b_thread_vec; - - static_for<0, KPack, 1>{}([&](auto ik) { - a_thread_vec.template AsType()(ik) = - a_thread_buf[Number{}]; - b_thread_vec.template AsType()(ik) = - b_thread_dequant_bufs[I1][Number{}]; - }); - - using mfma_input_type = - typename vector_type::type; - - constexpr index_t c_offset = - c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0)); - - xdlops_gemm.Run(a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf.GetVectorTypeReference(Number{})); - }); - }); - }); - // Let's leak last MFMA block to epilogue region, cover the potential lds-shuffle - // latency - // __builtin_amdgcn_sched_barrier(0); - } - else - { - static_for<0, MRepeat, 1>{}([&](auto m0) { - static_for<0, NRepeat, 1>{}([&](auto n0) { - static_for<0, KRepeat, 1>{}([&](auto k0) { - vector_type a_thread_vec; - vector_type b_thread_vec; - - static_for<0, KPack, 1>{}([&](auto ik) { - a_thread_vec.template AsType()(ik) = - a_thread_buf[Number{}]; - b_thread_vec.template AsType()(ik) = - b_thread_dequant_bufs[I0][Number{}]; - }); - - using mfma_input_type = - typename vector_type::type; - - constexpr index_t c_offset = - c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0)); - - xdlops_gemm.Run(a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf.GetVectorTypeReference(Number{})); - }); - }); - }); - } - } - - protected: - // MRepeat MWave MLane KRepeat KLane KPack - // KRepeat -> MRepeat-> Mwave->KLane->MLane->KPack - static constexpr auto a_thread_desc_ = make_naive_tensor_descriptor_packed( - make_tuple(Number{}, I1, I1, Number{}, I1, Number{})); - - using AThreadCopy = ThreadwiseTensorSliceTransfer_v4, - Sequence<0, 1, 2, 3, 4, 5>, - 5, - A_K1, - A_K1>; - - AThreadCopy a_thread_copy_{Base::CalculateAThreadOriginDataIndex6D()}; - - static constexpr auto b_thread_desc_ = make_naive_tensor_descriptor_packed( - make_tuple(Number{}, I1, Number{}, Number{})); - - static constexpr BTileDesc b_block_desc_n0_n1_k0_k1; - - using Base::c_thread_desc_; - - using PassThrough = ck::tensor_operation::element_wise::PassThrough; - - using BThreadDequantCopy = ThreadwiseTensorSliceTransfer_StaticToStatic< - BDataType, - ComputeDataType, - decltype(b_block_desc_n0_n1_k0_k1), - decltype(b_block_desc_n0_n1_k0_k1), - tensor_operation::element_wise::PassThrough, - Sequence{}, I1, Number{}, Number{}>, - Sequence<1, 2, 0, 3>, - 3, - KPack>; - - const PassThrough b_element_op{}; - BThreadDequantCopy b_thread_dequant_copy_{b_element_op}; -}; - -} // namespace ck diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v3.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v3.hpp deleted file mode 100644 index 83534ec5af..0000000000 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_dequant_v3.hpp +++ /dev/null @@ -1,928 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. - -#pragma once - -#include "ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_base.hpp" - -namespace ck { - -// Compute optimized pipeline -// GlobalPrefetchStages: 2 -// LocalPreFillStages: 1 -// LocalPreFetchStages: 1 -// LocalSharedMemoryBuffer: 1 - -template -struct BlockwiseGemmXdlops_pipeline_bpreshuffle_bdequant_v3 -{ -}; - -template -struct BlockwiseGemmXdlops_pipeline_bpreshuffle_bdequant_v3 - : BlockwiseGemmXdlops_pipeline_base - -{ - using Base = BlockwiseGemmXdlops_pipeline_base; - using Base::A_K1; - using Base::B_K1; - using Base::I0; - using Base::I1; - using Base::I2; - using Base::KRepeat; - using Base::xdlops_gemm; - using typename Base::HotLoopInstList; - - using Base::a_block_desc_m0_m1_m2_k; - using Base::CalculateCThreadOriginDataIndex; - using Base::CalculateCThreadOriginDataIndex8D; - using Base::GetCBlockDescriptor_G_M0_N0_M1_N1_M2_M3_M4_N2; - using Base::GetCBlockDescriptor_M0_N0_M1_N1_M2_M3_M4_N2; - using Base::GetCBlockDescriptor_M0_N0_M1_N1_M2_N2_N3_N4; - using Base::GetCThreadBuffer; - using Base::GetCThreadDescriptor_G_M0_N0_M1_N1_M2_M3_M4_N2; - using Base::GetCThreadDescriptor_M0_N0_M1_N1_M2_M3_M4_N2; - using Base::GetCThreadDescriptor_M0_N0_M1_N1_M2_N2_N3_N4; - using Base::MakeCGridDescriptor_G_M0_N0_M1_N1_M2_M3_M4_N2; - using Base::MakeCGridDescriptor_M0_N0_M1_N1_M2_M3_M4_N2; - - using Base::AMmaKStride; - using Base::BMmaKStride; - - using Base::MWaves; - - static constexpr index_t PrefetchStages = 2; - static constexpr index_t PrefillStages = 1; - static constexpr index_t GlobalBufferNum = 1; - static constexpr index_t HotloopLocalBufSwitch = MRepeat % 2 == 0 ? 0 : 1; - - template - __host__ __device__ static constexpr auto MakeAGemmMmaTileDescriptor(const TileDesc_M0_M1_M2_K&) - { - constexpr index_t M0 = TileDesc_M0_M1_M2_K{}.GetLength(Number<0>{}); - constexpr index_t M1 = TileDesc_M0_M1_M2_K{}.GetLength(Number<1>{}); - constexpr index_t M2 = TileDesc_M0_M1_M2_K{}.GetLength(Number<2>{}); - constexpr index_t K2 = KPack; - constexpr index_t K1 = 64 / NPerXDL; - constexpr index_t K0 = KRepeat; - - return transform_tensor_descriptor( - TileDesc_M0_M1_M2_K{}, - make_tuple( - make_pass_through_transform(Number{}), - make_pass_through_transform(Number{}), - make_pass_through_transform(Number{}), - make_unmerge_transform(make_tuple(Number{}, Number{}, Number{}))), - make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}), - make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3, 4, 5>{})); - } - - static constexpr auto a_block_desc_m0_m1_m2_k0_k1_k2 = - MakeAGemmMmaTileDescriptor(a_block_desc_m0_m1_m2_k); - - __host__ __device__ static constexpr bool BlockHasHotloop(index_t num_loop) - { - return num_loop > PrefetchStages; - } - - __host__ __device__ static constexpr TailNumber BlockLoopTailNum(index_t num_loop) - { - return num_loop % 2 == 0 ? TailNumber::Even : TailNumber::Odd; - } - - template - __device__ static constexpr auto HotLoopScheduler(Stage stage) - { - constexpr auto num_ds_read_inst_a = HotLoopInstList::A_LDS_Read_Inst_Num; - constexpr auto num_ds_write_inst_a = HotLoopInstList::A_LDS_Write_Inst_Num; - constexpr auto num_buffer_load_inst_a = HotLoopInstList::A_Buffer_Load_Inst_Num; - constexpr auto num_buffer_load_inst_b = MWaves * HotLoopInstList::B_Buffer_Load_Inst_Num; - - constexpr auto num_mfma = HotLoopInstList::C_MFMA_Inst_Num; - - constexpr auto staged_num_ds_read_inst_a = ck::math::integer_divide_ceil(num_ds_read_inst_a,MRepeat); - constexpr auto staged_num_mfma = ck::math::integer_divide_ceil(num_mfma , MRepeat); - - constexpr auto staged_num_mfma_per_ds_read_a = ck::math::integer_divide_ceil(staged_num_mfma , staged_num_ds_read_inst_a); - - if constexpr(stage.value == 0) - { - constexpr auto staged_num_buffer_load_b_per_ds_read_a = ck::math::integer_divide_ceil( - num_buffer_load_inst_b , staged_num_ds_read_inst_a); - constexpr auto staged_num_mfma_per_buffer_load_b =ck::math::integer_divide_ceil( - staged_num_mfma , num_buffer_load_inst_b); - // B global - static_for<0, staged_num_ds_read_inst_a, 1>{}([&](auto i_inst) { - ignore = i_inst; - - static_for<0, staged_num_buffer_load_b_per_ds_read_a - 1, 1>{}([&](auto ibuf_inst) { - ignore = ibuf_inst; - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_buffer_load_b, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read - }); - - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_buffer_load_b - 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read - }); - - __builtin_amdgcn_sched_barrier(0); - } - else if constexpr(stage.value == 1) - { - constexpr auto staged_num_mfma_per_ds_write_a = - math::integer_divide_ceil(staged_num_mfma, num_ds_write_inst_a); - - constexpr auto stage_more_mfma = - staged_num_mfma - (staged_num_mfma_per_ds_write_a - 1) * num_ds_write_inst_a; - - // A local write - static_for<0, num_ds_write_inst_a, 1>{}([&](auto i_inst) { - if constexpr(i_inst.value < stage_more_mfma) - { - if(i_inst.value < staged_num_ds_read_inst_a) - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_ds_write_a - 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS Write - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - } - else - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_ds_write_a, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS Write - } - } - else - { - if(i_inst.value < staged_num_ds_read_inst_a) - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_ds_write_a - 2, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS Write - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - } - else - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_ds_write_a - 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS Write - } - } - }); - - __builtin_amdgcn_sched_barrier(0); - } - else if constexpr(stage.value == 2) - { - constexpr auto staged_num_mfma_per_buffer_load_a = - math::integer_divide_ceil(staged_num_mfma, num_buffer_load_inst_a); - - constexpr auto stage_more_mfma = - staged_num_mfma - (staged_num_mfma_per_buffer_load_a - 1) * num_buffer_load_inst_a; - - // A global - static_for<0, num_buffer_load_inst_a, 1>{}([&](auto i_inst) { - if constexpr(i_inst.value < stage_more_mfma) - { - if(i_inst.value < staged_num_ds_read_inst_a) - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_buffer_load_a - 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - } - else - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_buffer_load_a, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read - } - } - else - { - if(i_inst.value < staged_num_ds_read_inst_a) - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_buffer_load_a - 2, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - } - else - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_buffer_load_a - 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read - } - } - }); - - __builtin_amdgcn_sched_barrier(0); - } - else - { - // A local Read - static_for<0, staged_num_ds_read_inst_a, 1>{}([&](auto i_inst) { - ignore = i_inst; - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_ds_read_a, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - }); - - __builtin_amdgcn_sched_barrier(0); - } - } - - template - __device__ static constexpr auto EpilogueScheduler_1(Stage stage) - { - constexpr auto num_ds_read_inst_a = HotLoopInstList::A_LDS_Read_Inst_Num; - constexpr auto num_ds_write_inst_a = HotLoopInstList::A_LDS_Write_Inst_Num; - constexpr auto num_buffer_load_inst_b = MWaves * HotLoopInstList::B_Buffer_Load_Inst_Num; - - constexpr auto num_mfma = HotLoopInstList::C_MFMA_Inst_Num; - - constexpr auto staged_num_ds_read_inst_a = num_ds_read_inst_a / MRepeat; - constexpr auto staged_num_mfma = num_mfma / MRepeat; - - constexpr auto staged_num_mfma_per_ds_read_a = staged_num_mfma / staged_num_ds_read_inst_a; - - if constexpr(stage.value == 0) - { - constexpr auto staged_num_buffer_load_b_per_ds_read_a = - num_buffer_load_inst_b / staged_num_ds_read_inst_a; - constexpr auto staged_num_mfma_per_buffer_load_b = - staged_num_mfma / num_buffer_load_inst_b; - // B global - static_for<0, staged_num_ds_read_inst_a, 1>{}([&](auto i_inst) { - ignore = i_inst; - - static_for<0, staged_num_buffer_load_b_per_ds_read_a, 1>{}([&](auto ibuf_inst) { - ignore = ibuf_inst; - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_buffer_load_b, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read - }); - - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_buffer_load_b - 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read - }); - - __builtin_amdgcn_sched_barrier(0); - } - else if constexpr(stage.value == 1) - { -#if 0 - constexpr auto staged_num_ds_write_a_per_ds_read_a = - num_ds_write_inst_a / staged_num_ds_read_inst_a; - constexpr auto staged_num_mfma_per_ds_write_a = staged_num_mfma / num_ds_write_inst_a; - // A local write - static_for<0, staged_num_ds_read_inst_a, 1>{}([&](auto i_inst) { - ignore = i_inst; - - static_for<0, staged_num_ds_write_a_per_ds_read_a, 1>{}([&](auto idswrite_inst) { - ignore = idswrite_inst; - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_ds_write_a - 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS Write - }); - - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_ds_write_a_per_ds_read_a, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - }); -#elif 1 - constexpr auto staged_num_mfma_per_ds_write_a = - math::integer_divide_ceil(staged_num_mfma, num_ds_write_inst_a); - - constexpr auto stage_more_mfma = - staged_num_mfma - (staged_num_mfma_per_ds_write_a - 1) * num_ds_write_inst_a; - - // A local write - static_for<0, num_ds_write_inst_a, 1>{}([&](auto i_inst) { - if constexpr(i_inst.value < stage_more_mfma) - { - if(i_inst.value < staged_num_ds_read_inst_a) - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_ds_write_a - 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS Write - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - } - else - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_ds_write_a, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS Write - } - } - else - { - if(i_inst.value < staged_num_ds_read_inst_a) - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_ds_write_a - 2, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS Write - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - } - else - { - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_ds_write_a - 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS Write - } - } - }); -#endif - __builtin_amdgcn_sched_barrier(0); - } - else - { - // A local Read - static_for<0, staged_num_ds_read_inst_a, 1>{}([&](auto i_inst) { - ignore = i_inst; - __builtin_amdgcn_sched_group_barrier( - 0x008, staged_num_mfma_per_ds_read_a, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - }); - - __builtin_amdgcn_sched_barrier(0); - } - } - - __device__ static constexpr auto EpilogueScheduler_2() - { - constexpr auto num_ds_read_inst_a = HotLoopInstList::A_LDS_Read_Inst_Num; - - constexpr auto num_mfma = HotLoopInstList::C_MFMA_Inst_Num; - - constexpr auto staged_num_ds_read_inst_a = num_ds_read_inst_a / MRepeat; - constexpr auto staged_num_mfma = num_mfma / MRepeat; - - constexpr auto staged_num_mfma_per_ds_read_a = staged_num_mfma / staged_num_ds_read_inst_a; - - // A local Read - static_for<0, staged_num_ds_read_inst_a, 1>{}([&](auto i_inst) { - ignore = i_inst; - __builtin_amdgcn_sched_group_barrier(0x008, staged_num_mfma_per_ds_read_a, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x100, 1, 0); // DS read - }); - - __builtin_amdgcn_sched_barrier(0); - } - - template - __device__ void Run(const AGridDesc& a_grid_desc, - const ABlockDesc& a_block_desc, - ABlockTransfer& a_blockwise_copy, - const AGridBuffer& a_grid_buf, - ABlockBuffer& a_block_buf, - const ABlockTransferStep& a_block_copy_step, - const BGridDesc& b_grid_desc, - BBlockTransfer& b_blockwise_copy, - const BGridBuffer& b_grid_buf, - BBlockBuffer& b_block_buf, - const BBlockTransferStep& b_block_copy_step, - CThreadBuffer& c_thread_buf, - index_t num_loop) const - { - ignore = b_block_buf; - __builtin_amdgcn_sched_barrier(0); - auto a_thread_buf = make_static_buffer( - a_thread_desc_.GetElementSpaceSize()); - auto b_thread_buf = make_static_buffer( - b_thread_desc_.GetElementSpaceSize()); - auto b_thread_dequant_buf = make_static_buffer( - b_thread_desc_.GetElementSpaceSize()); - - StaticallyIndexedArray{}> b_thread_bufs; - StaticallyIndexedArray{}> b_thread_dequant_bufs; - constexpr auto b_block_origin_idx = make_tuple(I0, I0, I0, I0); - - // Global prefetch A1 B1 - b_blockwise_copy.Run(b_grid_desc, - b_grid_buf, - b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(I0)); - b_blockwise_copy.MoveSrcSliceWindow(b_grid_desc, b_block_copy_step); - - a_blockwise_copy.RunRead(a_grid_desc, a_grid_buf); - a_blockwise_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step); - __builtin_amdgcn_sched_barrier(0); - - // // Local prefill A1 - a_blockwise_copy.RunWrite(a_block_desc, a_block_buf.At(I0)); - - // // Global prefetch A2 - a_blockwise_copy.RunRead(a_grid_desc, a_grid_buf); - a_blockwise_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step); - - // Local prefetch A1 - block_sync_lds(); - static_for<0, KRepeat, 1>{}([&](auto k0) { - a_thread_copy_.Run(a_block_desc_m0_m1_m2_k0_k1_k2, - make_tuple(I0, I0, I0, k0, I0, I0), - a_block_buf.At(I0), - a_thread_desc_, - make_tuple(I0, I0, I0, k0, I0, I0), - a_thread_buf); - }); - // B VGPR->VGPR dequant - b_thread_dequant_copy_.Run(b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(I0), - b_thread_desc_, - make_tuple(I0, I0, I0, I0), - b_thread_dequant_bufs(I0)); - - // Initialize C - c_thread_buf.Clear(); - - __builtin_amdgcn_sched_barrier(0); - - // main body - if constexpr(HasMainLoop) - { - index_t i = 0; - do - { - auto LoopFunc = [&](auto mfma_reg_buf, auto local_read_buf) { - static_for<0, MRepeat, 1>{}([&](auto m0) { - if constexpr(m0.value == 0) - { - b_blockwise_copy.Run(b_grid_desc, - b_grid_buf, - b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(local_read_buf)); - b_blockwise_copy.MoveSrcSliceWindow(b_grid_desc, b_block_copy_step); - } - else if constexpr(m0.value == 1) - { - a_blockwise_copy.RunWrite(a_block_desc, a_block_buf.At(local_read_buf)); - } - else if constexpr(m0.value == 2) - { - a_blockwise_copy.RunRead(a_grid_desc, a_grid_buf); - a_blockwise_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step); - } - - static_for<0, KRepeat, 1>{}([&](auto k0) { - static_for<0, NRepeat, 1>{}([&](auto n0) { - vector_type a_thread_vec; - vector_type b_thread_vec; - - static_for<0, KPack, 1>{}([&](auto ik) { - a_thread_vec.template AsType()(ik) = - a_thread_buf[Number{}]; - b_thread_vec.template AsType()(ik) = - b_thread_dequant_bufs[mfma_reg_buf] - [Number{}]; - }); - - using mfma_input_type = - typename vector_type::type; - - constexpr index_t c_offset = - c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0)); - - xdlops_gemm.Run( - a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf.GetVectorTypeReference(Number{})); - }); - }); - - if constexpr(m0.value == MRepeat - 1) - { - block_sync_lds(); - - static_for<0, KRepeat, 1>{}([&](auto k0) { - a_thread_copy_.Run( - a_block_desc_m0_m1_m2_k0_k1_k2, - make_tuple(Number<(m0 + 1) % MRepeat>{}, I0, I0, k0, I0, I0), - a_block_buf.At(local_read_buf), - a_thread_desc_, - make_tuple( - Number<(m0 + 1 + HotloopLocalBufSwitch * mfma_reg_buf) % - 2>{}, - I0, - I0, - k0, - I0, - I0), - a_thread_buf); - }); - // B VGPR->VGPR dequant - b_thread_dequant_copy_.Run(b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(local_read_buf), - b_thread_desc_, - make_tuple(I0, I0, I0, I0), - b_thread_dequant_bufs(local_read_buf)); - } - else - { - static_for<0, KRepeat, 1>{}([&](auto k0) { - a_thread_copy_.Run( - a_block_desc_m0_m1_m2_k0_k1_k2, - make_tuple(Number<(m0 + 1) % MRepeat>{}, I0, I0, k0, I0, I0), - a_block_buf.At(mfma_reg_buf), - a_thread_desc_, - make_tuple( - Number<(m0 + 1 + HotloopLocalBufSwitch * mfma_reg_buf) % - 2>{}, - I0, - I0, - k0, - I0, - I0), - a_thread_buf); - }); - // B VGPR->VGPR dequant - b_thread_dequant_copy_.Run(b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(mfma_reg_buf), - b_thread_desc_, - make_tuple(I0, I0, I0, I0), - b_thread_dequant_bufs(mfma_reg_buf)); - } - - HotLoopScheduler(m0); - }); - }; - - LoopFunc(I0, I1); - LoopFunc(I1, I0); - - i += 2; - } while(i < (num_loop - 2)); - } - // tail - if constexpr(TailNum == TailNumber::Even) - { - static_for<0, MRepeat, 1>{}([&](auto m0) { - if constexpr(m0.value == 0) - { - b_blockwise_copy.Run(b_grid_desc, - b_grid_buf, - b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(I1)); - } - else if constexpr(m0.value == MRepeat - 1) - { - a_blockwise_copy.RunWrite(a_block_desc, a_block_buf.At(I1)); - } - - static_for<0, KRepeat, 1>{}([&](auto k0) { - static_for<0, NRepeat, 1>{}([&](auto n0) { - vector_type a_thread_vec; - vector_type b_thread_vec; - - static_for<0, KPack, 1>{}([&](auto ik) { - a_thread_vec.template AsType()(ik) = - a_thread_buf[Number{}]; - b_thread_vec.template AsType()(ik) = - b_thread_dequant_bufs[I0][Number{}]; - }); - - using mfma_input_type = - typename vector_type::type; - - constexpr index_t c_offset = - c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0)); - - xdlops_gemm.Run(a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf.GetVectorTypeReference(Number{})); - }); - }); - - if constexpr(m0.value == MRepeat - 1) - { - block_sync_lds(); - - static_for<0, KRepeat, 1>{}([&](auto k0) { - a_thread_copy_.Run( - a_block_desc_m0_m1_m2_k0_k1_k2, - make_tuple(Number<(m0 + 1) % MRepeat>{}, I0, I0, k0, I0, I0), - a_block_buf.At(I1), - a_thread_desc_, - make_tuple(Number<(m0 + 1) % 2>{}, I0, I0, k0, I0, I0), - a_thread_buf); - }); - // B VGPR->VGPR dequant - b_thread_dequant_copy_.Run(b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(I1), - b_thread_desc_, - make_tuple(I0, I0, I0, I0), - b_thread_dequant_bufs(I1)); - } - else - { - static_for<0, KRepeat, 1>{}([&](auto k0) { - a_thread_copy_.Run( - a_block_desc_m0_m1_m2_k0_k1_k2, - make_tuple(Number<(m0 + 1) % MRepeat>{}, I0, I0, k0, I0, I0), - a_block_buf.At(I0), - a_thread_desc_, - make_tuple(Number<(m0 + 1) % 2>{}, I0, I0, k0, I0, I0), - a_thread_buf); - }); - // B VGPR->VGPR dequant - b_thread_dequant_copy_.Run(b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(I0), - b_thread_desc_, - make_tuple(I0, I0, I0, I0), - b_thread_dequant_bufs(I0)); - } - - EpilogueScheduler_1(m0); - }); - - static_for<0, MRepeat, 1>{}([&](auto m0) { - static_for<0, KRepeat, 1>{}([&](auto k0) { - static_for<0, NRepeat, 1>{}([&](auto n0) { - vector_type a_thread_vec; - vector_type b_thread_vec; - - static_for<0, KPack, 1>{}([&](auto ik) { - a_thread_vec.template AsType()(ik) = - a_thread_buf[Number{}]; - b_thread_vec.template AsType()(ik) = - b_thread_dequant_bufs[I1][Number{}]; - }); - - using mfma_input_type = - typename vector_type::type; - - constexpr index_t c_offset = - c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0)); - - xdlops_gemm.Run(a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf.GetVectorTypeReference(Number{})); - }); - }); - - if constexpr(m0.value != (MRepeat - 1)) - { - static_for<0, KRepeat, 1>{}([&](auto k0) { - a_thread_copy_.Run( - a_block_desc_m0_m1_m2_k0_k1_k2, - make_tuple(Number{}, I0, I0, k0, I0, I0), - a_block_buf.At(I1), - a_thread_desc_, - make_tuple( - Number<(m0 + 1 + HotloopLocalBufSwitch) % 2>{}, I0, I0, k0, I0, I0), - a_thread_buf); - }); - // B VGPR->VGPR dequant - b_thread_dequant_copy_.Run(b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(I1), - b_thread_desc_, - make_tuple(I0, I0, I0, I0), - b_thread_dequant_bufs(I1)); - - EpilogueScheduler_2(); - } - }); - // Let's leak last MFMA block to epilogue region, cover the potential lds-shuffle - // latency - // __builtin_amdgcn_sched_barrier(0); - } - else - { - static_for<0, MRepeat, 1>{}([&](auto m0) { - static_for<0, KRepeat, 1>{}([&](auto k0) { - static_for<0, NRepeat, 1>{}([&](auto n0) { - vector_type a_thread_vec; - vector_type b_thread_vec; - - static_for<0, KPack, 1>{}([&](auto ik) { - a_thread_vec.template AsType()(ik) = - a_thread_buf[Number{}]; - b_thread_vec.template AsType()(ik) = - b_thread_dequant_bufs[I0][Number{}]; - }); - - using mfma_input_type = - typename vector_type::type; - - constexpr index_t c_offset = - c_thread_desc_.CalculateOffset(make_tuple(m0, n0, 0)); - - xdlops_gemm.Run(a_thread_vec.template AsType(), - b_thread_vec.template AsType(), - c_thread_buf.GetVectorTypeReference(Number{})); - }); - }); - - if constexpr(m0.value != (MRepeat - 1)) - { - static_for<0, KRepeat, 1>{}([&](auto k0) { - a_thread_copy_.Run(a_block_desc_m0_m1_m2_k0_k1_k2, - make_tuple(Number{}, I0, I0, k0, I0, I0), - a_block_buf.At(I0), - a_thread_desc_, - make_tuple(Number<(m0 + 1) % 2>{}, I0, I0, k0, I0, I0), - a_thread_buf); - }); - // B VGPR->VGPR dequant - b_thread_dequant_copy_.Run(b_block_desc_n0_n1_k0_k1, - b_block_origin_idx, - b_thread_bufs(I0), - b_thread_desc_, - make_tuple(I0, I0, I0, I0), - b_thread_dequant_bufs(I0)); - - EpilogueScheduler_2(); - } - }); - } - } - - protected: - // MRepeat MWave MLane KRepeat KLane KPack - // KRepeat -> MRepeat-> Mwave->KLane->MLane->KPack - // Reduce the vgpr usage here. - static constexpr auto a_thread_desc_ = make_naive_tensor_descriptor_packed( - make_tuple(I2, I1, I1, Number{}, I1, Number{})); - - using AThreadCopy = ThreadwiseTensorSliceTransfer_v4, - Sequence<0, 1, 2, 3, 4, 5>, - 5, - A_K1, - A_K1>; - - AThreadCopy a_thread_copy_{Base::CalculateAThreadOriginDataIndex6D()}; - - static constexpr auto b_thread_desc_ = make_naive_tensor_descriptor_packed( - make_tuple(Number{}, I1, Number{}, Number{})); - - static constexpr BTileDesc b_block_desc_n0_n1_k0_k1; - - using Base::c_thread_desc_; - - using PassThrough = ck::tensor_operation::element_wise::PassThrough; - - using BThreadDequantCopy = ThreadwiseTensorSliceTransfer_StaticToStatic< - BDataType, - ComputeDataType, - decltype(b_block_desc_n0_n1_k0_k1), - decltype(b_block_desc_n0_n1_k0_k1), - tensor_operation::element_wise::PassThrough, - Sequence{}, I1, Number{}, Number{}>, - Sequence<1, 2, 0, 3>, - 3, - KPack>; - - const PassThrough b_element_op{}; - BThreadDequantCopy b_thread_dequant_copy_{b_element_op}; -}; - -} // namespace ck diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_selector.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_selector.hpp index 526c4216fa..0e647a10cb 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_selector.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_selector.hpp @@ -35,8 +35,6 @@ template ::value) { return BlockwiseGemmXdlops_pipeline_bpreshuffle_v1{}; - } - else - { - return BlockwiseGemmXdlops_pipeline_bpreshuffle_bdequant_v1< - BlkGemmPipeSche, - BlockSize, - ADataType, - BDataType, - ComputeDataType, - AccDataType, - ATileDesc, - BTileDesc, - AMmaTileDesc, - BMmaTileDesc, - ABlockTransferSrcScalarPerVector, - BBlockTransferSrcScalarPerVector, - MPerBlock, - NPerBlock, - KPerBlock, - MPerXDL, - NPerXDL, - MRepeat, - NRepeat, - KPack>{}; - } } else if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v2) { @@ -110,8 +83,6 @@ constexpr auto BlockGemmBPreshufflePipeline_Selector() else if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v3) { static_assert(MRepeat >= 4, "MRepeat should at least be 4 in BlockGemmPipelineVersion::v3"); - if constexpr(std::is_same::value) - { return BlockwiseGemmXdlops_pipeline_bpreshuffle_v3{}; - } - else - { - return BlockwiseGemmXdlops_pipeline_bpreshuffle_bdequant_v3< - BlkGemmPipeSche, - BlockSize, - ADataType, - BDataType, - ComputeDataType, - AccDataType, - ATileDesc, - BTileDesc, - AMmaTileDesc, - BMmaTileDesc, - ABlockTransferSrcScalarPerVector, - BBlockTransferSrcScalarPerVector, - MPerBlock, - NPerBlock, - KPerBlock, - MPerXDL, - NPerXDL, - MRepeat, - NRepeat, - KPack>{}; - } } else { diff --git a/include/ck/utility/amd_ck_fp8.hpp b/include/ck/utility/amd_ck_fp8.hpp index 42b784d303..0593a24bd3 100644 --- a/include/ck/utility/amd_ck_fp8.hpp +++ b/include/ck/utility/amd_ck_fp8.hpp @@ -6,6 +6,7 @@ #include "ck/ck.hpp" #include "ck/utility/enable_if.hpp" #include "ck/utility/random_gen.hpp" +#include "ck/utility/functional.hpp" #include "ck/utility/type.hpp" #ifdef CK_USE_FNUZ_FP8 @@ -193,10 +194,10 @@ __host__ __device__ static inline T cast_from_f8(fp8_storage_t x) } } - typename std::conditional< + typename ck::conditional_t< sizeof(T) == 2, unsigned short int, - typename std::conditional::type>::type + typename ck::conditional_t> retval; if constexpr(we == 5 && is_half && !is_fnuz) @@ -539,10 +540,10 @@ __host__ __device__ static inline fp8_storage_t cast_to_f8(T _x, unsigned int rn constexpr int mfmt = (sizeof(T) == 8) ? 52 : ((sizeof(T) == 4) ? 23 : 10); - using T_bitwise = typename std::conditional< + using T_bitwise = typename ck::conditional_t< sizeof(T) == 2, unsigned short int, - typename std::conditional::type>::type; + typename ck::conditional_t>; T_bitwise x_bitwise = bit_cast(_x); unsigned long long x{x_bitwise}; diff --git a/include/ck/utility/amd_inline_asm.hpp b/include/ck/utility/amd_inline_asm.hpp index de59f200f0..113f3af4ae 100644 --- a/include/ck/utility/amd_inline_asm.hpp +++ b/include/ck/utility/amd_inline_asm.hpp @@ -11,13 +11,6 @@ namespace ck { -inline __device__ int amd_assembly_and_b32(int a, int b) -{ - int c; - asm volatile("v_and_b32 %0, %1, %2" : "=v"(c) : "v"(a), "v"(b)); - return c; -} - inline __device__ int amd_assembly_and_or_b32(int a, int b, int d) { int c; @@ -39,54 +32,6 @@ inline __device__ half2_t amd_assembly_pk_add_f16(half2_t a, half2_t b) return c; } -inline __device__ float amd_assemble_cvt_f32_i4(int b) -{ - float a; - asm volatile("v_cvt_off_f32_i4 %0, %1" : "=v"(a) : "v"(b)); - return a; -} - -inline __device__ f8x4_t amd_assembly_cvt_f8_to_f32(float b0, float b1, float b2, float b3) -{ - f8x4_t a; - asm volatile("v_cvt_pk_fp8_f32 %0, %1, %2\n" - "v_cvt_pk_fp8_f32 %0, %3, %4, op_sel:[0, 0, 1]\n" - : "=v"(a) - : "v"(b0), "v"(b1), "v"(b2), "v"(b3)); - return a; -} - -inline __device__ f8x8_t amd_assembly_i4_to_fp8x8(int a) -{ - uint32_t i4x8 = static_cast(a); - uint32_t fp8x4_0; - uint32_t fp8x4_1; - float tmp_0, tmp_1, tmp_2; - - asm volatile("v_cvt_off_f32_i4 %[v_tmp_0], %[v_src]\n" - "v_cvt_off_f32_i4 %[v_tmp_1], %[v_src], src0_sel:BYTE_2\n" - "v_cvt_pk_fp8_f32 %[v_dst_0], %[v_tmp_0], %[v_tmp_1]\n" - "v_cvt_off_f32_i4 %[v_tmp_0], %[v_src], src0_sel:BYTE_1\n" - "v_cvt_off_f32_i4 %[v_tmp_1], %[v_src], src0_sel:BYTE_3\n" - "v_cvt_pk_fp8_f32 %[v_dst_1], %[v_tmp_0], %[v_tmp_1]\n" - "v_lshrrev_b32 %[v_tmp_2], 4, %[v_src]\n" - "v_cvt_off_f32_i4 %[v_tmp_0], %[v_tmp_2]\n" - "v_cvt_off_f32_i4 %[v_tmp_1], %[v_tmp_2], src0_sel:BYTE_2\n" - "v_cvt_pk_fp8_f32 %[v_dst_0], %[v_tmp_0], %[v_tmp_1], op_sel:[0, 0, 1]\n" - "v_cvt_off_f32_i4 %[v_tmp_0], %[v_tmp_2], src0_sel:BYTE_1\n" - "v_cvt_off_f32_i4 %[v_tmp_1], %[v_tmp_2], src0_sel:BYTE_3\n" - "v_cvt_pk_fp8_f32 %[v_dst_1], %[v_tmp_0], %[v_tmp_1], op_sel:[0, 0, 1]\n" - : [v_tmp_0] "+v"(tmp_0), - [v_tmp_1] "+v"(tmp_1), - [v_tmp_2] "+v"(tmp_2), - [v_dst_0] "+v"(fp8x4_0), - [v_dst_1] "+v"(fp8x4_1), - [v_src] "+v"(i4x8) - :); - - return bit_cast(((static_cast(fp8x4_1) << 32) | fp8x4_0)); -} - // c0 += inner_product(a, b0) // c1 += inner_product(a, b1) __device__ void amd_assembly_outer_product_1x2(float a, float b0, float b1, float& c0, float& c1) diff --git a/include/ck/utility/data_type.hpp b/include/ck/utility/data_type.hpp index 116804bf1c..a0d29e5a0f 100644 --- a/include/ck/utility/data_type.hpp +++ b/include/ck/utility/data_type.hpp @@ -19,7 +19,7 @@ using float_t = float; #endif // __HIPCC_RTC__ namespace ck { -#if !defined(__HIPCC_RTC__) || !defined(CK_CODE_GEN_RTC) +#if defined(__HIPCC_RTC__) || defined(CK_CODE_GEN_RTC) using byte = unsigned char; #else using std::byte; @@ -1191,15 +1191,11 @@ struct vector_type()>> StaticallyIndexedArray d8x4_; StaticallyIndexedArray d16x2_; StaticallyIndexedArray d32x1_; - } data_ = { .d32_ = {0} }; + } data_; - __attribute__((host)) __attribute__((device)) constexpr vector_type() { } - - __attribute__((host)) __attribute__((device)) constexpr vector_type(type v) { } + __host__ __device__ constexpr vector_type() : data_{type{0}} {} - // __host__ __device__ constexpr vector_type() : data_{type{0}} {} - - // __host__ __device__ constexpr vector_type(type v) : data_{v} {} + __host__ __device__ constexpr vector_type(type v) : data_{v} {} template __host__ __device__ constexpr const auto& AsType() const @@ -1809,7 +1805,7 @@ struct non_native_vector_base< // implementation for f6x16 and f6x32 template -struct non_native_vector_base> +struct non_native_vector_base> { using data_t = typename nnvb_data_t_selector::type; // select data_t based on declared base type diff --git a/include/ck/utility/mxf4_utils.hpp b/include/ck/utility/mxf4_utils.hpp index 15e693bd0d..757d3914e3 100644 --- a/include/ck/utility/mxf4_utils.hpp +++ b/include/ck/utility/mxf4_utils.hpp @@ -1,6 +1,7 @@ // SPDX-License-Identifier: MIT // Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. +#ifndef CK_CODE_GEN_RTC #pragma once #include "ck/utility/data_type.hpp" @@ -41,7 +42,7 @@ template <> __host__ __device__ inline float to_float(e8m0_bexp_t const scale, f4_t const data) { if(is_nan(scale, data)) - return std::numeric_limits::quiet_NaN(); + return NumericLimits::QuietNaN(); if(is_zero(scale, data)) return 0.0f; @@ -105,5 +106,5 @@ __host__ __device__ inline f4_t sat_convert_to_type_sr(float value, uint32 return res; } - } // namespace ck::utils +#endif diff --git a/include/ck/utility/mxf6_utils.hpp b/include/ck/utility/mxf6_utils.hpp index e3b37bedda..00b4f8e5d4 100644 --- a/include/ck/utility/mxf6_utils.hpp +++ b/include/ck/utility/mxf6_utils.hpp @@ -1,6 +1,7 @@ // SPDX-License-Identifier: MIT // Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. +#ifndef CK_CODE_GEN_RTC #pragma once #include "ck/utility/data_type.hpp" @@ -138,7 +139,7 @@ template <> __host__ __device__ inline float to_float(e8m0_bexp_t const scale, f6_t const data) { if(is_nan(scale, data)) - return std::numeric_limits::quiet_NaN(); + return NumericLimits::QuietNaN(); if(is_zero(scale, data)) return 0.0f; @@ -164,7 +165,7 @@ template <> __host__ __device__ inline float to_float(e8m0_bexp_t const scale, bf6_t const data) { if(is_nan(scale, data)) - return std::numeric_limits::quiet_NaN(); + return NumericLimits::QuietNaN(); if(is_zero(scale, data)) return 0.0f; @@ -307,7 +308,6 @@ __host__ __device__ inline bf6_t sat_convert_to_type_sr(float value, uint if(std::isnan(value)) return sign ? NumericUtils::data_max_negative_normal_mask : NumericUtils::data_max_positive_normal_mask; - if(std::abs(value) > NumericLimits::Max()) // covers inf case as well return sign ? NumericUtils::data_max_negative_normal_mask : NumericUtils::data_max_positive_normal_mask; @@ -321,5 +321,5 @@ __host__ __device__ inline bf6_t sat_convert_to_type_sr(float value, uint return res; } - } // namespace ck::utils +#endif diff --git a/include/ck/utility/mxfp_utils.hpp b/include/ck/utility/mxfp_utils.hpp index e23836c87f..947d64b705 100644 --- a/include/ck/utility/mxfp_utils.hpp +++ b/include/ck/utility/mxfp_utils.hpp @@ -3,6 +3,11 @@ #pragma once +#include "ck/utility/data_type.hpp" + +#ifdef CK_CODE_GEN_RTC +#define UINT_MAX 4294967295 +#endif namespace ck::utils { union cvt @@ -380,5 +385,4 @@ inline T convert_to_type_sr(float value, uint32_t seed) auto val = sign | biased_exp << NumericUtils::mant | mant; return val; } - } // namespace ck::utils diff --git a/include/ck/utility/type_convert.hpp b/include/ck/utility/type_convert.hpp index cf862ae640..69d1631ae3 100644 --- a/include/ck/utility/type_convert.hpp +++ b/include/ck/utility/type_convert.hpp @@ -706,7 +706,7 @@ inline __host__ __device__ half_t type_convert(bf8_fnuz_t x) return utils::cast_from_f8(x); #endif } - +#ifndef CK_CODE_GEN_RTC // convert fp32 to fp4 with rounding to nearest even inline __host__ __device__ f4_t f4_convert_rne(float x, float scale = 1.0f) { @@ -927,7 +927,11 @@ inline __host__ __device__ f4x32_t f4_convert_rne(float32_t x, float scale = 1.0 inline __host__ __device__ f4_t f4_convert_sr(float x, float scale = 1.0f) { constexpr int seed = 1254739; - uint32_t rng = prand_generator(reinterpret_cast(&x), x); +#ifndef CK_CODE_GEN_RTC + uint32_t rng = prand_generator(reinterpret_cast(&x), x); +#else + uint32_t rng = prand_generator(reinterpret_cast(&x), x); +#endif #if defined(__gfx950__) union { @@ -952,7 +956,11 @@ inline __host__ __device__ f4_t f4_convert_sr(float x, float scale = 1.0f) inline __host__ __device__ f4x2_t f4_convert_sr(float2_t x, float scale = 1.0f) { constexpr int seed = 1254739; - uint32_t rng = prand_generator(reinterpret_cast(&x), x[0]); +#ifndef CK_CODE_GEN_RTC + uint32_t rng = prand_generator(reinterpret_cast(&x), x[0]); +#else + uint32_t rng = prand_generator(reinterpret_cast(&x), x[0]); +#endif #if defined(__gfx950__) union { @@ -978,7 +986,11 @@ inline __host__ __device__ f4x2_t f4_convert_sr(float2_t x, float scale = 1.0f) inline __host__ __device__ f4x32_t f4_convert_sr(float32_t x, float scale = 1.0f) { constexpr int seed = 1254739; - uint32_t rng = prand_generator(reinterpret_cast(&x), x[0]); +#ifndef CK_CODE_GEN_RTC + uint32_t rng = prand_generator(reinterpret_cast(&x), x[0]); +#else + uint32_t rng = prand_generator(reinterpret_cast(&x), x[0]); +#endif #if defined(__gfx950__) union { @@ -1544,7 +1556,11 @@ inline __host__ __device__ f6x32_t f6_convert_rne(float32_t x, float scale = 1.0 inline __host__ __device__ f6_t f6_convert_sr(float x, float scale = 1.0f) { constexpr int seed = 1254739; - uint32_t rng = prand_generator(reinterpret_cast(&x), x); +#ifndef CK_CODE_GEN_RTC + uint32_t rng = prand_generator(reinterpret_cast(&x), x); +#else + uint32_t rng = prand_generator(reinterpret_cast(&x), x); +#endif #if defined(__gfx950__) union { @@ -1584,8 +1600,13 @@ inline __host__ __device__ f6x32_t f6_convert_sr(float32_t x, float scale = 1.0f float32_t float_vector; float float_array[32]; } float_values{x}; +#ifndef CK_CODE_GEN_RTC uint32_t rng = prand_generator(reinterpret_cast(&x), float_values.float_array[0]); +#else + uint32_t rng = + prand_generator(reinterpret_cast(&x), float_values.float_array[0]); +#endif #if defined(__gfx950__) return __builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f32(x, rng, scale); #else @@ -1803,7 +1824,11 @@ inline __host__ __device__ bf6x32_t bf6_convert_rne(float32_t x, float scale = 1 inline __host__ __device__ bf6_t bf6_convert_sr(float x, float scale = 1.0f) { constexpr int seed = 1254739; - uint32_t rng = prand_generator(reinterpret_cast(&x), x); +#ifndef CK_CODE_GEN_RTC + uint32_t rng = prand_generator(reinterpret_cast(&x), x); +#else + uint32_t rng = prand_generator(reinterpret_cast(&x), x); +#endif #if defined(__gfx950__) union { @@ -1845,8 +1870,13 @@ inline __host__ __device__ bf6x32_t bf6_convert_sr(float32_t x, float scale = 1. float32_t float_vector; float float_array[32]; } float_values{x}; +#ifndef CK_CODE_GEN_RTC uint32_t rng = prand_generator(reinterpret_cast(&x), float_values.float_array[0]); +#else + uint32_t rng = + prand_generator(reinterpret_cast(&x), float_values.float_array[0]); +#endif #if defined(__gfx950__) return __builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_f32(x, rng, scale); #else @@ -1978,7 +2008,7 @@ inline __host__ __device__ float32_t type_convert(bf6x32_t return out.float_vector; #endif } - +#endif #if !defined(__HIPCC_RTC__) || !defined(CK_CODE_GEN_RTC) template inline __host__ __device__ void array_convert(std::array& y, diff --git a/include/ck_tile/core/config.hpp b/include/ck_tile/core/config.hpp index 090b2bf797..b767f4b707 100644 --- a/include/ck_tile/core/config.hpp +++ b/include/ck_tile/core/config.hpp @@ -51,7 +51,7 @@ CK_TILE_DECLARE_ENV_VAR_BOOL(CK_TILE_LOGGING) // implementing the "memory address space" attribute // https://llvm.org/docs/AMDGPUUsage.html#amdgpu-address-spaces-table -#ifdef __HIPCC_ +#ifdef __HIPCC__ #define CK_TILE_GENERIC_ADDR __attribute__((address_space(0))) #define CK_TILE_GLOBAL_ADDR __attribute__((address_space(1))) #define CK_TILE_LDS_ADDR __attribute__((address_space(3))) diff --git a/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp b/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp index 0388798575..a1410d1f4f 100644 --- a/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp +++ b/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp @@ -581,11 +581,6 @@ struct MoeSortingKernel { int e_start = cumsum[tid]; int e_end = cumsum[tid + 1]; - index_t *p_sorted_expert_cnts = p_total_tokens_post_pad + 1 + tid; - if (tid == 0) { - p_sorted_expert_cnts[0] = 0; - } - p_sorted_expert_cnts[1] = unit_size_mdiv.div(e_end); for(int i = e_start; i < e_end; i += unit_size_mdiv.divisor) { p_sorted_expert_ids[unit_size_mdiv.div(i)] = tid; @@ -877,6 +872,7 @@ struct MoeSortingKernel } __syncthreads(); } + for(int i_e = tid; i_e < num_experts; i_e += block_size) { int e_start = smem_cumsum(i_e); @@ -891,6 +887,7 @@ struct MoeSortingKernel else return i_e; }(); + smem_cumdup(i_e) = e_start; // duplicate cumsum for later use if constexpr(Problem::SkipExpertsWithZeroTokens) { @@ -903,6 +900,7 @@ struct MoeSortingKernel if(local_expert_mask[i_e] == 0) continue; } + for(int i = e_start; i < e_end; i += unit_size_mdiv.divisor) { p_sorted_expert_ids[unit_size_mdiv.div(i)] = expert_id; @@ -986,18 +984,11 @@ struct MoeSortingKernel __syncthreads(); } - if (tid == 0) { - //temp hack ptr for expert tile cnt - p_total_tokens_post_pad[1] = 0; - } // add the skip number for(int eid = tid; eid < num_experts; eid += block_size) { - //temp hack ptr for expert tile cnt - index_t *p_sorted_expert_cnts = p_total_tokens_post_pad + 1 + eid; int e_start = smem_cumsum(eid); int e_end = smem_cumdup(eid + 1); - p_sorted_expert_cnts[1] = unit_size_mdiv.div(e_end); if constexpr(Problem::SkipExpertsWithZeroTokens) { if(e_start == e_end) // skip zero token expert @@ -1681,8 +1672,6 @@ struct MoeSortingMultiPhaseKernel_P2 if(position < kargs.num_experts) { - index_t *p_sorted_expert_cnts = p_total_tokens_post_pad + 1 + position;//temp mock for p_sorted_expert_cnts, fixme:felix - p_sorted_expert_cnts[0] = out_0; p_expert_cumsum[position] = out_0 * kargs.unit_size_mdiv.divisor; } @@ -1711,7 +1700,6 @@ struct MoeSortingMultiPhaseKernel_P2 { auto total_tokens_post_pad = prev_cumsum_a * kargs.unit_size_mdiv.divisor; p_total_tokens_post_pad[0] = total_tokens_post_pad; - p_total_tokens_post_pad[kargs.num_experts+1] = prev_cumsum_a; //temp mock for p_sorted_expert_cnts, fixme:felix p_expert_cumsum[kargs.num_experts] = total_tokens_post_pad; } } diff --git a/library/include/ck/library/tensor_operation_instance/gpu/gemm_ab_scale.hpp b/library/include/ck/library/tensor_operation_instance/gpu/gemm_ab_scale.hpp index 3fa82ae53a..7553d5e76e 100644 --- a/library/include/ck/library/tensor_operation_instance/gpu/gemm_ab_scale.hpp +++ b/library/include/ck/library/tensor_operation_instance/gpu/gemm_ab_scale.hpp @@ -17,7 +17,7 @@ namespace tensor_operation { namespace device { namespace instance { #if(defined(CK_ENABLE_BF16) || defined(CK_ENABLE_FP8)) -void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_default_instances( +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_default_instances( std::vector, @@ -28,14 +28,14 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_default_ins F32, Tuple<>, BF16, - 1, + 128, 128, 128, PassThrough, PassThrough, PassThrough>>>& instances); -void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_kpadding_instances( +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_kpadding_instances( std::vector, @@ -46,14 +46,14 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_kpadding_in F32, Tuple<>, BF16, - 1, + 128, 128, 128, PassThrough, PassThrough, PassThrough>>>& instances); -void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_default_instances( +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnpadding_instances( std::vector, @@ -64,14 +64,14 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_default_i F32, Tuple<>, BF16, - 1, + 128, 128, 128, PassThrough, PassThrough, PassThrough>>>& instances); -void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_kpadding_instances( +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnkpadding_instances( std::vector, @@ -82,7 +82,61 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_kpadding_ F32, Tuple<>, BF16, - 1, + 128, + 128, + 128, + PassThrough, + PassThrough, + PassThrough>>>& instances); + +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_default_instances( + std::vector, + Row, + F8, + F32, + F8, + F32, + Tuple<>, + BF16, + 128, + 128, + 128, + PassThrough, + PassThrough, + PassThrough>>>& instances); + +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_kpadding_instances( + std::vector, + Row, + F8, + F32, + F8, + F32, + Tuple<>, + BF16, + 128, + 128, + 128, + PassThrough, + PassThrough, + PassThrough>>>& instances); + +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_mnkpadding_instances( + std::vector, + Row, + F8, + F32, + F8, + F32, + Tuple<>, + BF16, + 128, 128, 128, PassThrough, @@ -109,7 +163,7 @@ struct DeviceOperationInstanceFactory, CDataType, - 1, + 128, 128, 128, ck::tensor_operation::element_wise::PassThrough, @@ -126,7 +180,7 @@ struct DeviceOperationInstanceFactory, CDataType, - 1, + 128, 128, 128, ck::tensor_operation::element_wise::PassThrough, @@ -144,14 +198,20 @@ struct DeviceOperationInstanceFactory && is_same_v && is_same_v) { - add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_default_instances( + add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_default_instances( op_ptrs); - add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_kpadding_instances( + add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_kpadding_instances( + op_ptrs); + add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnpadding_instances( + op_ptrs); + add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnkpadding_instances( op_ptrs); - add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_default_instances( + add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_default_instances( op_ptrs); - add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_kpadding_instances( + add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_kpadding_instances( + op_ptrs); + add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_mnkpadding_instances( op_ptrs); } } diff --git a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/CMakeLists.txt b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/CMakeLists.txt index d572862884..aab1c4e86e 100644 --- a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/CMakeLists.txt +++ b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/CMakeLists.txt @@ -4,13 +4,16 @@ set(GEMM_AB_SCALE_INSTANCES) list(APPEND GEMM_AB_SCALE_INSTANCES device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_default_instance.cpp device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_kpadding_instance.cpp + device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnpadding_instance.cpp + device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnkpadding_instance.cpp device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_default_instance.cpp device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_kpadding_instance.cpp + device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_mnkpadding_instance.cpp ) set_source_files_properties(device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_default_instance.cpp PROPERTIES COMPILE_OPTIONS ";-mllvm;-greedy-reverse-local-assignment=1") set_source_files_properties(device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_kpadding_instance.cpp PROPERTIES COMPILE_OPTIONS ";-mllvm;-greedy-reverse-local-assignment=1") -set_source_files_properties(device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_default_instance.cpp PROPERTIES COMPILE_OPTIONS ";-mllvm;-greedy-reverse-local-assignment=1") -set_source_files_properties(device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_kpadding_instance.cpp PROPERTIES COMPILE_OPTIONS ";-mllvm;-greedy-reverse-local-assignment=1") +set_source_files_properties(device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnpadding_instance.cpp PROPERTIES COMPILE_OPTIONS ";-mllvm;-greedy-reverse-local-assignment=1") +set_source_files_properties(device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnkpadding_instance.cpp PROPERTIES COMPILE_OPTIONS ";-mllvm;-greedy-reverse-local-assignment=1") add_instance_library(device_gemm_ab_scale_instance ${GEMM_AB_SCALE_INSTANCES}) diff --git a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128.hpp b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128.hpp index eba9cfcb7c..3a7df8d974 100644 --- a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128.hpp +++ b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128.hpp @@ -34,50 +34,49 @@ static constexpr auto GemmMNKPadding = GemmSpecialization::MNKPadding; static constexpr auto Intrawave = BlockGemmPipelineScheduler::Intrawave; template -using device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_instances = std::tuple< +using device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_instances = std::tuple< // clang-format off - //################################| ALayout| BLayout| DsLayout| ELayout| AData| BData| DsData| EData| AccData| Cshuffle| A| B| C| GEMM| Block| Scale| Scale| Scale| MPer| NPer| KPer| AK1| BK1|MPer| NPer| MXdl| NXdl| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockLds| BBlockTransfer| BBlockTransfer| BBlockTransfer| BlockTransfer| BBlockTransfer| BBlockTransfer| BBlockLds| CShuffle| CShuffle| CBlockTransferClusterLengths| CBlockTransfer| Block-wiseGemm| Block-wiseGemm| - //################################| | | | | Type| Type| Type| Type| Type| Type| Elementwise| Elementwise| Elementwise|Specialization| Size| Block| Block| Block| Block| Block| Block| | | XDL| XDL| Per| Per| ThreadCluster| ThreadCluster| SrcAccessOrder| SrcVectorDim| SrcScalar| DstScalar| AddExtraM| ThreadCluster| ThreadCluster| SrcAccessOrder| SrcVectorDim| SrcScalar| DstScalar| AddExtraN| MXdlPerWave| NXdlPerWave| _MBlock_MXdlPerWave_MWaveMPerXdl| ScalarPerVector| Pipeline| Pipeline| - //################################| | | | | | | | | | | Operation| Operation| Operation| | | M| N| K| | | | | | | | Wave| Wave| Lengths_K0_M_K1| ArrangeOrder| | | PerVector| PerVector_K1| | Lengths_K0_N_K1| ArrangeOrder| | | PerVector| PerVector_K1| | PerShuffle| PerShuffle| _NBlock_NXdlPerWave_NWaveNPerXdl| _NWaveNPerXdl| Scheduler| Verision| - //################################| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | + //################################| ALayout| BLayout| DsLayout| ELayout|AData| BData| DsData| EData| AccData| Cshuffle| A| B| C| GEMM| Block| Scale| Scale| Scale| MPer| NPer| KPer| AK1| BK1|MPer| NPer| MXdl| NXdl| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockLds| BBlockTransfer| BBlockTransfer| BBlockTransfer| BlockTransfer| BBlockTransfer| BBlockTransfer| BBlockLds| CShuffle| CShuffle| CBlockTransferClusterLengths| CBlockTransfer| Block-wiseGemm| Block-wiseGemm| + //################################| | | | | Type| Type| Type| Type| Type| Type| Elementwise| Elementwise| Elementwise|Specialization| Size| Block| Block| Block| Block| Block| Block| | | XDL| XDL| Per| Per| ThreadCluster| ThreadCluster| SrcAccessOrder| SrcVectorDim| SrcScalar| DstScalar| AddExtraM| ThreadCluster| ThreadCluster| SrcAccessOrder| SrcVectorDim| SrcScalar| DstScalar| AddExtraN| MXdlPerWave| NXdlPerWave| _MBlock_MXdlPerWave_MWaveMPerXdl| ScalarPerVector| Pipeline| Pipeline| + //################################| | | | | | | | | | | Operation| Operation| Operation| | | M| N| K| | | | | | | | Wave| Wave| Lengths_K0_M_K1| ArrangeOrder| | | PerVector| PerVector_K1| | Lengths_K0_N_K1| ArrangeOrder| | | PerVector| PerVector_K1| | PerShuffle| PerShuffle| _NBlock_NXdlPerWave_NWaveNPerXdl| _NWaveNPerXdl| Scheduler| Verision| + //################################| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | // Compute friendly - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8, F32, F8, F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 128, 128, 128, 16, 16, 32, 32, 2, 2, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v3, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8, F32, F8, F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 128, 64, 128, 16, 16, 32, 32, 2, 1, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v3, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8, F32, F8, F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 64, 128, 128, 16, 16, 32, 32, 1, 2, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v3, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8, F32, F8, F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 64, 64, 128, 16, 16, 32, 32, 1, 1, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v3, F8> + // Spill in current compiler + // DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8, F8, Tuple, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 224, 256, 128, 16, 16, 16, 16, 7, 8, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 2, S<1, 32, 1, 8>, S<8, 8, 1>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v3, F8>, + // DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8, F8, Tuple, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 256, 224, 128, 16, 16, 16, 16, 8, 7, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 2, 1, S<1, 64, 1, 4>, S<8, 8, 1>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v3, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8, F32, F8, F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 128, 128, 128, 128, 128, 128, 16, 16, 32, 32, 2, 2, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8, 8, 1>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v3, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8, F32, F8, F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 128, 128, 128, 128, 64, 128, 16, 16, 32, 32, 2, 1, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8, 8, 1>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v3, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8, F32, F8, F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 128, 128, 128, 64, 128, 128, 16, 16, 32, 32, 1, 2, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8, 8, 1>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v3, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8, F32, F8, F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 128, 128, 128, 64, 64, 128, 16, 16, 32, 32, 1, 1, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8, 8, 1>, BlockGemmPipelineScheduler::Intrawave, BlockGemmPipelineVersion::v3, F8> // clang-format on >; template -using device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_instances = std::tuple< +using device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_instances = std::tuple< // clang-format off - //################################| ALayout| BLayout| DsLayout| ELayout|AData | BData| DsData| EData| AccData| Cshuffle| A| B| C| GEMM| Block| Scale| Scale| Scale| MPer| NPer| KPer| AK1| BK1|MPer| NPer| MXdl| NXdl| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockLds| BBlockTransfer| BBlockTransfer| BBlockTransfer| BlockTransfer| BBlockTransfer| BBlockTransfer| BBlockLds| CShuffle| CShuffle| CBlockTransferClusterLengths| CBlockTransfer| Block-wiseGemm| Block-wiseGemm| - //################################| | | | | Type | Type| Type| Type| Type| Type| Elementwise| Elementwise| Elementwise|Specialization| Size| Block| Block| Block| Block| Block| Block| | | XDL| XDL| Per| Per| ThreadCluster| ThreadCluster| SrcAccessOrder| SrcVectorDim| SrcScalar| DstScalar| AddExtraM| ThreadCluster| ThreadCluster| SrcAccessOrder| SrcVectorDim| SrcScalar| DstScalar| AddExtraN| MXdlPerWave| NXdlPerWave| _MBlock_MXdlPerWave_MWaveMPerXdl| ScalarPerVector| Pipeline| Pipeline| - //################################| | | | | | | | | | | Operation| Operation| Operation| | | M| N| K| | | | | | | | Wave| Wave| Lengths_K0_M_K1| ArrangeOrder| | | PerVector| PerVector_K1| | Lengths_K0_N_K1| ArrangeOrder| | | PerVector| PerVector_K1| | PerShuffle| PerShuffle| _NBlock_NXdlPerWave_NWaveNPerXdl| _NWaveNPerXdl| Scheduler| Verision| - //################################| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | + //################################| ALayout| BLayout| DsLayout| ELayout|AData| BData| DsData| EData| AccData| Cshuffle| A| B| C| GEMM| Block| Scale| Scale| Scale| MPer| NPer| KPer| AK1| BK1|MPer| NPer| MXdl| NXdl| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockLds| BBlockTransfer| BBlockTransfer| BBlockTransfer| BlockTransfer| BBlockTransfer| BBlockTransfer| BBlockLds| CShuffle| CShuffle| CBlockTransferClusterLengths| CBlockTransfer| Block-wiseGemm| Block-wiseGemm| + //################################| | | | | Type| Type| Type| Type| Type| Type| Elementwise| Elementwise| Elementwise|Specialization| Size| Block| Block| Block| Block| Block| Block| | | XDL| XDL| Per| Per| ThreadCluster| ThreadCluster| SrcAccessOrder| SrcVectorDim| SrcScalar| DstScalar| AddExtraM| ThreadCluster| ThreadCluster| SrcAccessOrder| SrcVectorDim| SrcScalar| DstScalar| AddExtraN| MXdlPerWave| NXdlPerWave| _MBlock_MXdlPerWave_MWaveMPerXdl| ScalarPerVector| Pipeline| Pipeline| + //################################| | | | | | | | | | | Operation| Operation| Operation| | | M| N| K| | | | | | | | Wave| Wave| Lengths_K0_M_K1| ArrangeOrder| | | PerVector| PerVector_K1| | Lengths_K0_N_K1| ArrangeOrder| | | PerVector| PerVector_K1| | PerShuffle| PerShuffle| _NBlock_NXdlPerWave_NWaveNPerXdl| _NWaveNPerXdl| Scheduler| Verision| + //################################| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | - // Memory friendly - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 16, 256, 128, 8, 16, 16, 16, 1, 4, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 2, S<1, 16, 1, 16>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 16, 128, 128, 8, 16, 16, 16, 1, 2, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 2, S<1, 16, 1, 16>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 16, 64, 128, 8, 16, 16, 16, 1, 1, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 16>, S<4>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 16, 128, 256, 16, 16, 16, 16, 1, 2, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 2, S<1, 16, 1, 16>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 16, 64, 256, 16, 16, 16, 16, 1, 1, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 16>, S<4>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 32, 256, 128, 16, 16, 32, 32, 1, 2, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 32, 128, 128, 16, 16, 32, 32, 1, 1, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 32, 64, 128, 16, 16, 16, 16, 2, 1, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 2, 1, S<1, 32, 1, 8>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 32, 128, 256, 16, 16, 32, 32, 1, 1, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 32, 64, 256, 16, 16, 16, 16, 2, 1, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 2, 1, S<1, 32, 1, 8>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 64, 256, 128, 16, 16, 32, 32, 2, 2, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 64, 128, 128, 16, 16, 32, 32, 2, 1, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 64, 64, 128, 16, 16, 32, 32, 1, 1, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 64, 128, 256, 16, 16, 32, 32, 2, 1, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, - DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 256, 1, 128, 128, 64, 64, 256, 16, 16, 32, 32, 1, 1, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<16, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 32, 1, 8>, S<8>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8> + // Latency friendly + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 32, 16, 128, 16, 16, 16, 16, 1, 1, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<2, 2, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 64, 128, 128, 128, 16, 16, 128, 16, 16, 16, 16, 1, 1, S<8, 8, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 8, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 4>, S<4, 4, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 16, 32, 128, 16, 16, 16, 16, 1, 1, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<4, 4, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v1, F8>, + // Memory friendly + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 128, 32, 128, 16, 16, 32, 32, 2, 1, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<4, 4, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 128, 16, 128, 16, 16, 16, 16, 4, 1, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<2, 2, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 64, 32, 128, 16, 16, 32, 32, 1, 1, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<4, 4, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 64, 16, 128, 16, 16, 16, 16, 2, 1, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<2, 2, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 32, 16, 128, 16, 16, 16, 16, 1, 1, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<2, 2, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 64, 128, 128, 128, 16, 16, 64, 16, 16, 16, 16, 1, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 4>, S<4, 4, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 64, 128, 128, 128, 16, 16, 128, 16, 16, 16, 16, 1, 1, S<8, 8, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 8, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 4>, S<4, 4, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 16, 32, 128, 16, 16, 16, 16, 1, 1, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<4, 4, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 16, 64, 128, 16, 16, 16, 16, 1, 2, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<4, 4, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 32, 64, 128, 16, 16, 32, 32, 1, 1, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<8, 8, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 16, 128, 128, 16, 16, 16, 16, 1, 4, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<4, 4, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8>, + DeviceGemmMultiD_ABScale_Xdl_CShuffle_V3< Row, Col, Tuple<>, Row, F8,F32, F8,F32, Tuple<>, BF16, F32, F32, PassThrough, PassThrough, PassThrough, GemmSpec, 128, 128, 128, 128, 32, 128, 128, 16, 16, 32, 32, 1, 2, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, S<8, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 16, 16, 0, 1, 1, S<1, 16, 1, 8>, S<8, 8, 1>, BlkGemmPipeSched, BlockGemmPipelineVersion::v2, F8> // clang-format on >; } // namespace instance diff --git a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_default_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_default_instance.cpp index aebffc01f2..ab83c7eb3e 100644 --- a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_default_instance.cpp +++ b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_default_instance.cpp @@ -8,7 +8,7 @@ namespace tensor_operation { namespace device { namespace instance { -void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_default_instances( +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_default_instances( std::vector, @@ -19,7 +19,7 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_default_ins F32, Tuple<>, BF16, - 1, + 128, 128, 128, PassThrough, @@ -28,7 +28,7 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_default_ins { add_device_operation_instances( instances, - device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_instances{}); + device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_instances{}); } } // namespace instance diff --git a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_kpadding_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_kpadding_instance.cpp index 31fffae080..dfb1bb6e2d 100644 --- a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_kpadding_instance.cpp +++ b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_kpadding_instance.cpp @@ -8,7 +8,7 @@ namespace tensor_operation { namespace device { namespace instance { -void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_kpadding_instances( +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_kpadding_instances( std::vector, @@ -19,7 +19,7 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_kpadding_in F32, Tuple<>, BF16, - 1, + 128, 128, 128, PassThrough, @@ -28,7 +28,7 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_kpadding_in { add_device_operation_instances( instances, - device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_comp_instances{}); + device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_instances{}); } } // namespace instance diff --git a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnkpadding_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnkpadding_instance.cpp new file mode 100644 index 0000000000..d2d3ebe81e --- /dev/null +++ b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnkpadding_instance.cpp @@ -0,0 +1,37 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include "device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { +namespace instance { + +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnkpadding_instances( + std::vector, + Row, + F8, + F32, + F8, + F32, + Tuple<>, + BF16, + 128, + 128, + 128, + PassThrough, + PassThrough, + PassThrough>>>& instances) +{ + add_device_operation_instances( + instances, + device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_instances{}); +} + +} // namespace instance +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnpadding_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnpadding_instance.cpp new file mode 100644 index 0000000000..f6ce77a751 --- /dev/null +++ b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnpadding_instance.cpp @@ -0,0 +1,37 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include "device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { +namespace instance { + +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_mnpadding_instances( + std::vector, + Row, + F8, + F32, + F8, + F32, + Tuple<>, + BF16, + 128, + 128, + 128, + PassThrough, + PassThrough, + PassThrough>>>& instances) +{ + add_device_operation_instances( + instances, + device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_comp_instances{}); +} + +} // namespace instance +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_default_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_default_instance.cpp index 569911e3de..e2205ad728 100644 --- a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_default_instance.cpp +++ b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_default_instance.cpp @@ -8,7 +8,7 @@ namespace tensor_operation { namespace device { namespace instance { -void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_default_instances( +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_default_instances( std::vector, @@ -19,7 +19,7 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_default_i F32, Tuple<>, BF16, - 1, + 128, 128, 128, PassThrough, @@ -28,8 +28,8 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_default_i { add_device_operation_instances( instances, - device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_instances{}); + device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_instances{}); } } // namespace instance diff --git a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_kpadding_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_kpadding_instance.cpp index d1e5b6b535..5c0a6eb00d 100644 --- a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_kpadding_instance.cpp +++ b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_kpadding_instance.cpp @@ -8,7 +8,7 @@ namespace tensor_operation { namespace device { namespace instance { -void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_kpadding_instances( +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_kpadding_instances( std::vector, @@ -19,7 +19,7 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_kpadding_ F32, Tuple<>, BF16, - 1, + 128, 128, 128, PassThrough, @@ -28,8 +28,8 @@ void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_v1_kpadding_ { add_device_operation_instances( instances, - device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_1_128_128_mem_instances{}); + device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_instances{}); } } // namespace instance diff --git a/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_mnkpadding_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_mnkpadding_instance.cpp new file mode 100644 index 0000000000..cc1a03b060 --- /dev/null +++ b/library/src/tensor_operation_instance/gpu/gemm_ab_scale/device_gemm_ab_scale_xdl_f8_f8_bf16/device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_mnkpadding_instance.cpp @@ -0,0 +1,38 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + +#include "device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128.hpp" + +namespace ck { +namespace tensor_operation { +namespace device { +namespace instance { + +void add_device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_v1_mnkpadding_instances( + std::vector, + Row, + F8, + F32, + F8, + F32, + Tuple<>, + BF16, + 128, + 128, + 128, + PassThrough, + PassThrough, + PassThrough>>>& instances) +{ + add_device_operation_instances( + instances, + device_gemm_ab_scale_xdl_f8_f8_bf16_mk_nk_mn_128_128_128_mem_instances{}); +} + +} // namespace instance +} // namespace device +} // namespace tensor_operation +} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p1_padding_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p1_padding_instance.cpp deleted file mode 100644 index 00e860e022..0000000000 --- a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p1_padding_instance.cpp +++ /dev/null @@ -1,33 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - -#include "device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn.hpp" - -namespace ck { -namespace tensor_operation { -namespace device { -namespace instance { - -void add_device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p1_padding_instances( - std::vector, - Row, - F8, - F8, - Tuple, - BF16, - PassThrough, - PassThrough, - MultiplyMultiply>>>& instances) -{ - add_device_operation_instances( - instances, - device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p1_instances< - GemmKPadding>{}); -} - -} // namespace instance -} // namespace device -} // namespace tensor_operation -} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p2_padding_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p2_padding_instance.cpp deleted file mode 100644 index 4ef710857d..0000000000 --- a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p2_padding_instance.cpp +++ /dev/null @@ -1,33 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - -#include "device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn.hpp" - -namespace ck { -namespace tensor_operation { -namespace device { -namespace instance { - -void add_device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p2_padding_instances( - std::vector, - Row, - F8, - F8, - Tuple, - BF16, - PassThrough, - PassThrough, - MultiplyMultiply>>>& instances) -{ - add_device_operation_instances( - instances, - device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p2_instances< - GemmKPadding>{}); -} - -} // namespace instance -} // namespace device -} // namespace tensor_operation -} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p3_padding_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p3_padding_instance.cpp deleted file mode 100644 index 6700722ae5..0000000000 --- a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p3_padding_instance.cpp +++ /dev/null @@ -1,33 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - -#include "device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn.hpp" - -namespace ck { -namespace tensor_operation { -namespace device { -namespace instance { - -void add_device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p3_padding_instances( - std::vector, - Row, - F8, - F8, - Tuple, - BF16, - PassThrough, - PassThrough, - MultiplyMultiply>>>& instances) -{ - add_device_operation_instances( - instances, - device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_bf16_mk_mfma_mn_p3_instances< - GemmKPadding>{}); -} - -} // namespace instance -} // namespace device -} // namespace tensor_operation -} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p1_padding_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p1_padding_instance.cpp deleted file mode 100644 index 8b48f3cba0..0000000000 --- a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p1_padding_instance.cpp +++ /dev/null @@ -1,33 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - -#include "device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn.hpp" - -namespace ck { -namespace tensor_operation { -namespace device { -namespace instance { - -void add_device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p1_padding_instances( - std::vector, - Row, - F8, - F8, - Tuple, - F16, - PassThrough, - PassThrough, - MultiplyMultiply>>>& instances) -{ - add_device_operation_instances( - instances, - device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p1_instances< - GemmKPadding>{}); -} - -} // namespace instance -} // namespace device -} // namespace tensor_operation -} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p2_padding_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p2_padding_instance.cpp deleted file mode 100644 index 579dba5a78..0000000000 --- a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p2_padding_instance.cpp +++ /dev/null @@ -1,33 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - -#include "device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn.hpp" - -namespace ck { -namespace tensor_operation { -namespace device { -namespace instance { - -void add_device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p2_padding_instances( - std::vector, - Row, - F8, - F8, - Tuple, - F16, - PassThrough, - PassThrough, - MultiplyMultiply>>>& instances) -{ - add_device_operation_instances( - instances, - device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p2_instances< - GemmKPadding>{}); -} - -} // namespace instance -} // namespace device -} // namespace tensor_operation -} // namespace ck diff --git a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p3_padding_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p3_padding_instance.cpp deleted file mode 100644 index ee7ce201af..0000000000 --- a/library/src/tensor_operation_instance/gpu/gemm_multiply_multiply_weight_preshuffle/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16/device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p3_padding_instance.cpp +++ /dev/null @@ -1,33 +0,0 @@ -// SPDX-License-Identifier: MIT -// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - -#include "device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn.hpp" - -namespace ck { -namespace tensor_operation { -namespace device { -namespace instance { - -void add_device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p3_padding_instances( - std::vector, - Row, - F8, - F8, - Tuple, - F16, - PassThrough, - PassThrough, - MultiplyMultiply>>>& instances) -{ - add_device_operation_instances( - instances, - device_gemm_multiply_multiply_weight_preshuffle_xdl_f8_f8_f16_mk_mfma_mn_p3_instances< - GemmKPadding>{}); -} - -} // namespace instance -} // namespace device -} // namespace tensor_operation -} // namespace ck diff --git a/profiler/src/profile_gemm_ab_scale.cpp b/profiler/src/profile_gemm_ab_scale.cpp index 3956038a30..56c8b5e7a1 100644 --- a/profiler/src/profile_gemm_ab_scale.cpp +++ b/profiler/src/profile_gemm_ab_scale.cpp @@ -32,7 +32,6 @@ enum struct GemmDataType enum struct ScaleBlockTile { Tile_128_128_128, // 0 - Tile_1_128_128, // 1 }; #define OP_NAME "gemm_ab_scale" @@ -50,8 +49,7 @@ int profile_gemm_ab_scale(int argc, char* argv[]) printf(" 1: A[m, k] * B[n, k] = C[m, n];\n"); printf(" 2: A[k, m] * B[k, n] = C[m, n];\n"); printf(" 3: A[k, m] * B[n, k] = C[m, n])\n"); - printf("arg4: scale block tile (0: ScaleBlockM/N/K = [128, 128, 128]; 1: ScaleBlockM/N/K = " - "[1, 128, 128];\n"); + printf("arg4: scale block tile (0: ScaleBlockM/N/K = [128, 128, 128];\n"); printf("arg5: verification (0: no; 1: yes)\n"); printf("arg6: initialization (0: no init; 1: integer value; 2: decimal value)\n"); printf("arg7: print tensor value (0: no; 1: yes)\n"); @@ -157,7 +155,7 @@ int profile_gemm_ab_scale(int argc, char* argv[]) }; if(data_type == GemmDataType::F8_F8_BF16 && layout == GemmMatrixLayout::MK_NK_MN && - scale_block_tile == ScaleBlockTile::Tile_1_128_128) + scale_block_tile == ScaleBlockTile::Tile_128_128_128) { return profile(F8{}, F32{}, @@ -166,7 +164,7 @@ int profile_gemm_ab_scale(int argc, char* argv[]) F8{}, F32{}, BF16{}, - ck::Number<1>{}, + ck::Number<128>{}, ck::Number<128>{}, ck::Number<128>{}, Row{},