ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-15 03:30:11 +00:00
Code Issues Packages Projects Releases Wiki Activity

1. copy bhalf v3 to half v3

Browse Source 
2. add multi config in half v2
3. sync block universal with half v3
This commit is contained in:
Qun Lin
2025-05-12 15:07:31 +08:00
parent b1ed92b131
commit efe964bd95
12 changed files with 215 additions and 82 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
example/01_gemm/CMakeLists.txt
View File

@@ -18,13 +18,17 @@ endif(USE_BITINT_EXTENSION_INT4)
add_custom_target(example_gemm_xdl)
add_example_executable(example_gemm_xdl_fp16 gemm_xdl_fp16.cpp)
add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16)
target_compile_options(example_gemm_xdl_fp16 PRIVATE -save-temps=obj -Wno-gnu-line-marker)
add_example_executable(example_gemm_xdl_fp16_v2 gemm_xdl_fp16_v2.cpp)
add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16_v2)
target_compile_options(example_gemm_xdl_fp16_v2 PRIVATE -save-temps=obj -Wno-gnu-line-marker)
add_example_executable(example_gemm_xdl_fp16_streamk_v3 gemm_xdl_fp16_streamk_v3.cpp)
add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16_streamk_v3)
add_example_executable(example_gemm_xdl_fp16_v3 gemm_xdl_fp16_v3.cpp)
target_compile_options(example_gemm_xdl_fp16_v3 PRIVATE -save-temps=obj -Wno-gnu-line-marker)
add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16_v3)
add_example_executable(example_gemm_xdl_fp8_v3 gemm_xdl_fp8_v3.cpp)
add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp8_v3)
@@ -38,6 +42,7 @@ add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16_fp8_streamk_v3)
add_example_executable(example_gemm_xdl_bf16_v3 gemm_xdl_bf16_v3.cpp)
add_example_dependencies(example_gemm_xdl example_gemm_xdl_bf16_v3)
target_compile_options(example_gemm_xdl_bf16_v3 PRIVATE -save-temps=obj -Wno-gnu-line-marker)
list(APPEND gpu_list gfx942 gfx950)
set(target 0)
@@ -60,6 +65,7 @@ add_example_dependencies(example_gemm_xdl example_gemm_xdl_skip_b_lds_fp16)
add_example_executable(example_gemm_xdl_bf16 gemm_xdl_bf16.cpp)
add_example_dependencies(example_gemm_xdl example_gemm_xdl_bf16)
target_compile_options(example_gemm_xdl_bf16 PRIVATE -save-temps=obj -Wno-gnu-line-marker)
add_example_executable(example_gemm_xdl_int8 gemm_xdl_int8.cpp)
add_example_dependencies(example_gemm_xdl example_gemm_xdl_int8)

78
example/01_gemm/gemm_xdl_fp16_v2.cpp
View File

@@ -59,4 +59,80 @@ using ReferenceGemmInstanceGPU = ck::tensor_operation::device::ReferenceGemm<ALa
#include "run_gemm_example.inc"
int main(int argc, char* argv[]) { return !run_gemm_example(argc, argv); }
struct GemmTypeConifg_RR
{
using ADataType_ = ck::half_t;
using BDataType_ = ck::half_t;
using CDataType_ = ck::half_t;
using ALayout_ = Row;
using BLayout_ = Row;
using CLayout_ = Row;
};
struct GemmTypeConifg_RC
{
using ADataType_ = ck::half_t;
using BDataType_ = ck::half_t;
using CDataType_ = ck::half_t;
using ALayout_ = Row;
using BLayout_ = Col;
using CLayout_ = Row;
};
using DeviceGemmInstance_0 =
ck::tensor_operation::device::DeviceGemm_Xdl_CShuffleV2<
ALayout, BLayout, CLayout,
F16, F16, F16, F32, F16,
PassThrough, PassThrough, PassThrough, GemmDefault,
2, 256,
256, 256,
32, 8, 4,
32, 32,
4, 4,
S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>,
2, 8, 8, 0,
S<8, 32, 1>, S<0, 2, 1>, S<0, 2, 1>,
1, 8, 4, 0,
1, 1, S<1, 32, 1, 8>, 8,
ck::LoopScheduler::Default, ck::PipelineVersion::v1>;
using DeviceGemmInstance_1 =
ck::tensor_operation::device::DeviceGemm_Xdl_CShuffleV2<
ALayout, BLayout, CLayout,
F16, F16, F16, F32, F16,
PassThrough, PassThrough, PassThrough, GemmDefault,
2, 256,
256, 256,
32, 8, 8,
32, 32,
4, 4,
S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>,
2, 8, 8, 0,
S<4, 32, 2>, S<0, 2, 1>, S<0, 2, 1>,
1, 8, 8, 0,
1, 1, S<1, 32, 1, 8>, 8,
ck::LoopScheduler::Default, ck::PipelineVersion::v1>;
using DeviceGemmInstance_2 =
ck::tensor_operation::device::DeviceGemm_Xdl_CShuffleV2<
ALayout, BLayout, CLayout,
F16, F16, F16, F32, F16,
PassThrough, PassThrough, PassThrough, GemmDefault,
2, 256,
256, 256,
32, 8, 8,
32, 32,
4, 4,
S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>,
2, 8, 8, 0,
S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>,
2, 8, 8, 0,
1, 1, S<1, 32, 1, 8>, 8,
ck::LoopScheduler::Default, ck::PipelineVersion::v1>;
int main(int argc, char* argv[])
{
run_gemm_example_with_instance<DeviceGemmInstance_0, GemmTypeConifg_RR>(argc, argv);
run_gemm_example_with_instance<DeviceGemmInstance_1, GemmTypeConifg_RR>(argc, argv);
run_gemm_example_with_instance<DeviceGemmInstance_2, GemmTypeConifg_RC>(argc, argv);
}

18
example/01_gemm/gemm_xdl_fp16_v3.cpp
View File

@@ -19,7 +19,7 @@ using AElementOp = PassThrough;
using BElementOp = PassThrough;
using CElementOp = PassThrough;
static constexpr auto GemmDefault = ck::tensor_operation::device::GemmSpecialization::MNPadding;
static constexpr auto GemmDefault = ck::tensor_operation::device::GemmSpecialization::Default;
// clang-format off
using DeviceGemmV2Instance =
@@ -27,17 +27,17 @@ using DeviceGemmV2Instance =
ALayout, BLayout, CLayout,
ADataType, BDataType, CDataType, AccDataType, CShuffleDataType,
PassThrough, PassThrough, PassThrough, GemmDefault,
64,
16, 16,
256, 8, 8,
256,
128, 128,
64, 8, 8,
16, 16,
1, 1,
S<32, 2, 1>, S<1, 0, 2>, S<1, 0, 2>,
4, 4,
S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>,
2, 8, 8, 0,
S<32, 2, 1>, S<1, 0, 2>, S<1, 0, 2>,
S<8, 32, 1>, S<1, 0, 2>, S<1, 0, 2>,
2, 8, 8, 0,
1, 1, S<1, 16, 1, 4>, 4,
ck::BlockGemmPipelineScheduler::Interwave, ck::BlockGemmPipelineVersion::v2>;
1, 2, S<1, 32, 1, 8>, 8,
ck::BlockGemmPipelineScheduler::Intrawave,ck::BlockGemmPipelineVersion::v3>;
// clang-format on
using ReferenceGemmInstance = ck::tensor_operation::host::

116
example/01_gemm/run_gemm_example.inc
View File

@@ -3,7 +3,7 @@
#pragma once
template <typename ProblemType>
template <typename DeviceGemm, typename GemmConfig, typename ProblemType>
bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
{
#if defined(BUILD_INT4_EXAMPLE) && defined(CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4)
@@ -19,6 +19,13 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
auto StrideB = problem_size.StrideB;
auto StrideC = problem_size.StrideC;
using ALayout_ = typename GemmConfig::ALayout_;
using BLayout_ = typename GemmConfig::BLayout_;
using CLayout_ = typename GemmConfig::CLayout_;
using ADataType_ = typename GemmConfig::ADataType_;
using BDataType_ = typename GemmConfig::BDataType_;
using CDataType_ = typename GemmConfig::CDataType_;
auto f_host_tensor_descriptor =
[](std::size_t row, std::size_t col, std::size_t stride, auto layout) {
if constexpr(std::is_same_v<decltype(layout), ck::tensor_layout::gemm::RowMajor>)
@@ -49,68 +56,68 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
return static_cast<std::size_t>(stride);
};
StrideA = f_get_default_stride(M, K, StrideA, ALayout{});
StrideB = f_get_default_stride(K, N, StrideB, BLayout{});
StrideC = f_get_default_stride(M, N, StrideC, CLayout{});
StrideA = f_get_default_stride(M, K, StrideA, ALayout_{});
StrideB = f_get_default_stride(K, N, StrideB, BLayout_{});
StrideC = f_get_default_stride(M, N, StrideC, CLayout_{});
Tensor<ADataType> a_m_k(f_host_tensor_descriptor(M, K, StrideA, ALayout{}));
Tensor<BDataType> b_k_n(f_host_tensor_descriptor(K, N, StrideB, BLayout{}));
Tensor<ADataType_> a_m_k(f_host_tensor_descriptor(M, K, StrideA, ALayout_{}));
Tensor<BDataType_> b_k_n(f_host_tensor_descriptor(K, N, StrideB, BLayout_{}));
switch(config.init_method)
{
case 0:
ck::utils::FillConstant<ADataType>{ck::type_convert<ADataType>(1.f)}(a_m_k);
ck::utils::FillConstant<BDataType>{ck::type_convert<BDataType>(1.f)}(b_k_n);
ck::utils::FillConstant<ADataType_>{ck::type_convert<ADataType_>(1.f)}(a_m_k);
ck::utils::FillConstant<BDataType_>{ck::type_convert<BDataType_>(1.f)}(b_k_n);
break;
case 1:
ck::utils::FillUniformDistributionIntegerValue<ADataType>{-5.f, 5.f}(a_m_k);
ck::utils::FillUniformDistributionIntegerValue<BDataType>{-5.f, 5.f}(b_k_n);
ck::utils::FillUniformDistributionIntegerValue<ADataType_>{-5.f, 5.f}(a_m_k);
ck::utils::FillUniformDistributionIntegerValue<BDataType_>{-5.f, 5.f}(b_k_n);
break;
case 2:
ck::utils::FillUniformDistribution<ADataType>{-1.f, 1.f}(a_m_k);
ck::utils::FillUniformDistribution<BDataType>{-1.f, 1.f}(b_k_n);
ck::utils::FillUniformDistribution<ADataType_>{-1.f, 1.f}(a_m_k);
ck::utils::FillUniformDistribution<BDataType_>{-1.f, 1.f}(b_k_n);
break;
case 3:
ck::utils::FillUniformDistributionIntegerValue<ADataType>{1.f, 1.f}(a_m_k);
ck::utils::FillUniformDistributionIntegerValue<BDataType>{-5.f, 5.f}(b_k_n);
ck::utils::FillUniformDistributionIntegerValue<ADataType_>{1.f, 1.f}(a_m_k);
ck::utils::FillUniformDistributionIntegerValue<BDataType_>{-5.f, 5.f}(b_k_n);
break;
case 4:
ck::utils::FillUniformDistributionIntegerValue<ADataType>{-5.f, 5.f}(a_m_k);
ck::utils::FillUniformDistributionIntegerValue<BDataType>{1.f, 1.f}(b_k_n);
ck::utils::FillUniformDistributionIntegerValue<ADataType_>{-5.f, 5.f}(a_m_k);
ck::utils::FillUniformDistributionIntegerValue<BDataType_>{1.f, 1.f}(b_k_n);
break;
case 5:
ck::utils::FillUniformDistributionIntegerValue<ADataType>{-2.f, 2.f}(a_m_k);
ck::utils::FillUniformDistributionIntegerValue<BDataType>{-2.f, 2.f}(b_k_n);
ck::utils::FillUniformDistributionIntegerValue<ADataType_>{-2.f, 2.f}(a_m_k);
ck::utils::FillUniformDistributionIntegerValue<BDataType_>{-2.f, 2.f}(b_k_n);
break;
default:
ck::utils::FillUniformDistribution<ADataType>{-0.1f, 0.1f}(a_m_k);
ck::utils::FillUniformDistribution<BDataType>{-0.1f, 0.1f}(b_k_n);
ck::utils::FillUniformDistribution<ADataType_>{-0.1f, 0.1f}(a_m_k);
ck::utils::FillUniformDistribution<BDataType_>{-0.1f, 0.1f}(b_k_n);
}
Tensor<CDataType> c_m_n_host_result(f_host_tensor_descriptor(M, N, StrideC, CLayout{}));
Tensor<CDataType> c_m_n_device_result(f_host_tensor_descriptor(M, N, StrideC, CLayout{}));
Tensor<CDataType> c_m_n_device_ref_result(f_host_tensor_descriptor(M, N, StrideC, CLayout{}));
Tensor<CDataType_> c_m_n_host_result(f_host_tensor_descriptor(M, N, StrideC, CLayout_{}));
Tensor<CDataType_> c_m_n_device_result(f_host_tensor_descriptor(M, N, StrideC, CLayout_{}));
Tensor<CDataType_> c_m_n_device_ref_result(f_host_tensor_descriptor(M, N, StrideC, CLayout_{}));
std::cout << "a_m_k: " << a_m_k.mDesc << std::endl;
std::cout << "b_k_n: " << b_k_n.mDesc << std::endl;
std::cout << "c_m_n: " << c_m_n_host_result.mDesc << std::endl;
#ifdef BUILD_INT4_EXAMPLE
DeviceMem a_m_k_device_buf(sizeof(KernelADataType) * a_m_k.mDesc.GetElementSpaceSize());
DeviceMem a_m_k_device_buf(sizeof(KernelADataType_) * a_m_k.mDesc.GetElementSpaceSize());
DeviceMem b_k_n_device_buf(sizeof(KernelBDataType) * b_k_n.mDesc.GetElementSpaceSize());
DeviceMem c_m_n_device_buf(sizeof(KernelCDataType) *
c_m_n_device_result.mDesc.GetElementSpaceSize());
const Tensor<KernelADataType> a_m_k_converted(a_m_k);
const Tensor<KernelADataType_> a_m_k_converted(a_m_k);
const Tensor<KernelBDataType> b_k_n_converted(b_k_n);
a_m_k_device_buf.ToDevice(a_m_k_converted.mData.data());
b_k_n_device_buf.ToDevice(b_k_n_converted.mData.data());
#else
DeviceMem a_m_k_device_buf(sizeof(ADataType) * a_m_k.mDesc.GetElementSpaceSize());
DeviceMem b_k_n_device_buf(sizeof(BDataType) * b_k_n.mDesc.GetElementSpaceSize());
DeviceMem c_m_n_device_buf(sizeof(CDataType) * c_m_n_device_result.mDesc.GetElementSpaceSize());
DeviceMem c_m_n_device_ref_buf(sizeof(CDataType) *
DeviceMem a_m_k_device_buf(sizeof(ADataType_) * a_m_k.mDesc.GetElementSpaceSize());
DeviceMem b_k_n_device_buf(sizeof(BDataType_) * b_k_n.mDesc.GetElementSpaceSize());
DeviceMem c_m_n_device_buf(sizeof(CDataType_) * c_m_n_device_result.mDesc.GetElementSpaceSize());
DeviceMem c_m_n_device_ref_buf(sizeof(CDataType_) *
c_m_n_device_ref_result.mDesc.GetElementSpaceSize());
a_m_k_device_buf.ToDevice(a_m_k.mData.data());
@@ -123,7 +130,7 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
auto c_element_op = CElementOp{};
// do GEMM
auto gemm = DeviceGemmInstance{};
auto gemm = DeviceGemm{};
auto invoker = gemm.MakeInvoker();
float ave_time = 0;
@@ -131,13 +138,13 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
{
auto argument = gemm.MakeArgument(
#ifdef BUILD_INT4_EXAMPLE
static_cast<KernelADataType*>(a_m_k_device_buf.GetDeviceBuffer()),
static_cast<KernelADataType_*>(a_m_k_device_buf.GetDeviceBuffer()),
static_cast<KernelBDataType*>(b_k_n_device_buf.GetDeviceBuffer()),
static_cast<KernelCDataType*>(c_m_n_device_buf.GetDeviceBuffer()),
#else
static_cast<ADataType*>(a_m_k_device_buf.GetDeviceBuffer()),
static_cast<BDataType*>(b_k_n_device_buf.GetDeviceBuffer()),
static_cast<CDataType*>(c_m_n_device_buf.GetDeviceBuffer()),
static_cast<ADataType_*>(a_m_k_device_buf.GetDeviceBuffer()),
static_cast<BDataType_*>(b_k_n_device_buf.GetDeviceBuffer()),
static_cast<CDataType_*>(c_m_n_device_buf.GetDeviceBuffer()),
#endif
M,
N,
@@ -169,7 +176,7 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
std::size_t flop = 2_uz * M * N * K;
std::size_t num_btype =
sizeof(ADataType) * M * K + sizeof(BDataType) * K * N + sizeof(CDataType) * M * N;
sizeof(ADataType_) * M * K + sizeof(BDataType_) * K * N + sizeof(CDataType_) * M * N;
float tflops = static_cast<float>(flop) / 1.E9 / ave_time;
@@ -193,11 +200,11 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
ref_invoker.Run(ref_argument);
#ifdef BUILD_INT4_EXAMPLE
Tensor<CDataType> c_m_n_device_result_converted(c_m_n_host_result.mDesc);
Tensor<CDataType_> c_m_n_device_result_converted(c_m_n_host_result.mDesc);
c_m_n_device_buf.FromDevice(c_m_n_device_result_converted.mData.data());
c_m_n_device_result = c_m_n_device_result_converted.CopyAsType<CDataType>();
c_m_n_device_result = c_m_n_device_result_converted.CopyAsType<CDataType_>();
return ck::utils::check_err(c_m_n_device_result_converted, c_m_n_host_result);
#else
@@ -206,8 +213,8 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
pass &= ck::utils::check_err(c_m_n_device_result,
c_m_n_host_result,
"Error: Incorrect results!",
get_rtol<CDataType>(),
get_atol<CDataType>());
get_rtol<CDataType_>(),
get_atol<CDataType_>());
#endif
}
@@ -218,9 +225,9 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
auto ref_invoker_gpu = ref_gemm_gpu.MakeInvoker();
auto ref_argument_gpu = ref_gemm_gpu.MakeArgument(
static_cast<ADataType*>(a_m_k_device_buf.GetDeviceBuffer()),
static_cast<BDataType*>(b_k_n_device_buf.GetDeviceBuffer()),
static_cast<CDataType*>(c_m_n_device_ref_buf.GetDeviceBuffer()),
static_cast<ADataType_*>(a_m_k_device_buf.GetDeviceBuffer()),
static_cast<BDataType_*>(b_k_n_device_buf.GetDeviceBuffer()),
static_cast<CDataType_*>(c_m_n_device_ref_buf.GetDeviceBuffer()),
M,
N,
K,
@@ -237,17 +244,36 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
pass &= ck::utils::check_err(c_m_n_device_result,
c_m_n_device_ref_result,
"Error: Incorrect results!",
get_rtol<CDataType>(),
get_atol<CDataType>());
get_rtol<CDataType_>(),
get_atol<CDataType_>());
}
return pass == true;
}
struct GemmConfigDefault
{
using ALayout_ = ALayout;
using BLayout_ = BLayout;
using CLayout_ = CLayout;
using ADataType_ = ADataType;
using BDataType_ = BDataType;
using CDataType_ = CDataType;
};
bool run_gemm_example(int argc, char* argv[])
{
ProblemSize problem_size;
ExecutionConfig config;
return !parse_cmd_args(argc, argv, problem_size, config) || run_gemm(problem_size, config);
return !parse_cmd_args(argc, argv, problem_size, config) || run_gemm<DeviceGemmInstance, GemmConfigDefault>(problem_size, config);
}
template<typename DeviceGemm, typename GemmConfig>
bool run_gemm_example_with_instance(int argc, char* argv[])
{
ProblemSize problem_size;
ExecutionConfig config;
return !parse_cmd_args(argc, argv, problem_size, config) || run_gemm<DeviceGemm, GemmConfig>(problem_size, config);
}

5
example/ck_tile/03_gemm/CMakeLists.txt
View File

@@ -7,3 +7,8 @@ endif()
list(APPEND EXAMPLE_GEMM_COMPILE_OPTIONS -mllvm -enable-noalias-to-md-conversion=0)
target_compile_options(tile_example_gemm_basic PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})
target_compile_options(tile_example_gemm_universal PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})
target_compile_options(tile_example_gemm_universal PRIVATE -save-temps=obj -Wno-gnu-line-marker)
target_compile_options(tile_example_gemm_basic PRIVATE -save-temps=obj -Wno-gnu-line-marker)
#target_compile_options(tile_example_gemm_basic PRIVATE ";-mllvm;--slp-threshold=-32")
#target_compile_options(tile_example_gemm_universal PRIVATE ";-mllvm;--slp-threshold=-32")

22
example/ck_tile/03_gemm/gemm_basic.cpp
View File

@@ -26,11 +26,11 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config&
constexpr bool kPadN = false;
constexpr bool kPadK = false;
constexpr int kBlockPerCu = 1;
constexpr int kBlockPerCu = 2;
// This part comes from the Codegen
constexpr ck_tile::index_t M_Tile = 256;
constexpr ck_tile::index_t N_Tile = 256;
constexpr ck_tile::index_t M_Tile = 128;
constexpr ck_tile::index_t N_Tile = 128;
constexpr ck_tile::index_t K_Tile = 64;
constexpr ck_tile::index_t M_Warp = 2;
@@ -144,16 +144,17 @@ int run_gemm_example_prec_type(std::string a_layout, std::string b_layout, int a
}
else
{
if(a_layout == "R" && b_layout == "R")
{
return run_gemm_example_with_layouts<APrecType, BPrecType, CPrecType>(
argc, argv, Row{}, Row{}, Row{});
}
else if(a_layout == "R" && b_layout == "C")
if(a_layout == "R" && b_layout == "C")
{
return run_gemm_example_with_layouts<APrecType, BPrecType, CPrecType>(
argc, argv, Row{}, Col{}, Row{});
}
#if 0
else if(a_layout == "R" && b_layout == "R")
{
return run_gemm_example_with_layouts<APrecType, BPrecType, CPrecType>(
argc, argv, Row{}, Row{}, Row{});
}
else if(a_layout == "C" && b_layout == "R")
{
return run_gemm_example_with_layouts<APrecType, BPrecType, CPrecType>(
@@ -164,6 +165,7 @@ int run_gemm_example_prec_type(std::string a_layout, std::string b_layout, int a
return run_gemm_example_with_layouts<APrecType, BPrecType, CPrecType>(
argc, argv, Col{}, Col{}, Row{});
}
#endif
else
{
throw std::runtime_error("Unsupported memory layout for the input matrices!");
@@ -185,6 +187,7 @@ int run_gemm_example(int argc, char* argv[])
{
return run_gemm_example_prec_type<ck_tile::half_t>(a_layout, b_layout, argc, argv);
}
#if 0
else if(data_type == "bf16")
{
return run_gemm_example_prec_type<ck_tile::bf16_t>(a_layout, b_layout, argc, argv);
@@ -207,6 +210,7 @@ int run_gemm_example(int argc, char* argv[])
return run_gemm_example_prec_type<ck_tile::half_t, ck_tile::pk_int4_t, ck_tile::half_t>(
a_layout, b_layout, argc, argv);
}
#endif
#endif
else
{

4
example/ck_tile/03_gemm/gemm_utils.hpp
View File

@@ -57,7 +57,7 @@ struct GemmConfig
// Compute friendly for Intrawave scheduler
static constexpr ck_tile::index_t M_Tile = 128;
static constexpr ck_tile::index_t N_Tile = 128;
static constexpr ck_tile::index_t K_Tile = 128;
static constexpr ck_tile::index_t K_Tile = 64;
static constexpr ck_tile::index_t M_Warp = 2;
static constexpr ck_tile::index_t N_Warp = 2;
@@ -96,7 +96,7 @@ struct GemmConfig
static constexpr bool TransposeC = false;
static constexpr bool UseStructuredSparsity = false;
static constexpr int kBlockPerCu = 1;
static constexpr int kBlockPerCu = 2;
static constexpr ck_tile::index_t TileParitionerGroupNum = 8;
static constexpr ck_tile::index_t TileParitionerM01 = 4;
};

30
example/ck_tile/03_gemm/universal_gemm.cpp
View File

@@ -41,10 +41,9 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config&
sequence<GemmConfig::M_Warp_Tile, GemmConfig::N_Warp_Tile, GemmConfig::K_Warp_Tile>,
GemmConfig::PermuteA,
GemmConfig::PermuteB>;
using TilePartitioner =
ck_tile::GemmSpatiallyLocalTilePartitioner<GemmShape,
GemmConfig::TileParitionerGroupNum,
GemmConfig::TileParitionerM01>;
using TilePartitioner = ck_tile::GemmSpatiallyLocalTilePartitioner<GemmShape,
GemmConfig::TileParitionerGroupNum,
GemmConfig::TileParitionerM01>;
using Traits = ck_tile::TileGemmTraits<GemmConfig::kPadM,
GemmConfig::kPadN,
@@ -141,6 +140,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config&
ck_tile::integral_constant<ck_tile::memory_operation_enum,
ck_tile::memory_operation_enum::set>{});
}
#if 0
else
{
Run(has_hot_loop_,
@@ -148,6 +148,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config&
ck_tile::integral_constant<ck_tile::memory_operation_enum,
ck_tile::memory_operation_enum::atomic_add>{});
}
#endif
};
if(has_hot_loop)
@@ -158,6 +159,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config&
RunSplitk(ck_tile::bool_constant<true>{},
ck_tile::integral_constant<ck_tile::TailNumber, ck_tile::TailNumber::Full>{});
}
#if 0
else if(tail_num == ck_tile::TailNumber::Odd)
{
RunSplitk(ck_tile::bool_constant<true>{},
@@ -168,6 +170,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config&
RunSplitk(ck_tile::bool_constant<true>{},
ck_tile::integral_constant<ck_tile::TailNumber, ck_tile::TailNumber::Even>{});
}
#endif
else
{
std::ostringstream err;
@@ -215,6 +218,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config&
}
else
{
#if 0
if(tail_num == ck_tile::TailNumber::Full)
{
RunSplitk(ck_tile::bool_constant<false>{},
@@ -231,6 +235,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config&
ck_tile::integral_constant<ck_tile::TailNumber, ck_tile::TailNumber::Even>{});
}
else
#endif
{
std::ostringstream err;
err << "Num K loop must be larger than number of prefetech stages."
@@ -271,16 +276,18 @@ int run_gemm_example_prec_type(std::string a_layout, std::string b_layout, int a
}
else
{
if(a_layout == "R" && b_layout == "R")
{
return run_gemm_example_with_layouts<APrecType, BPrecType, CPrecType>(
argc, argv, Row{}, Row{}, Row{});
}
else if(a_layout == "R" && b_layout == "C")
if(a_layout == "R" && b_layout == "C")
{
return run_gemm_example_with_layouts<APrecType, BPrecType, CPrecType>(
argc, argv, Row{}, Col{}, Row{});
}
#if 0
else if(a_layout == "R" && b_layout == "R")
{
return run_gemm_example_with_layouts<APrecType, BPrecType, CPrecType>(
argc, argv, Row{}, Row{}, Row{});
}
else if(a_layout == "C" && b_layout == "R")
{
return run_gemm_example_with_layouts<APrecType, BPrecType, CPrecType>(
@@ -291,6 +298,7 @@ int run_gemm_example_prec_type(std::string a_layout, std::string b_layout, int a
return run_gemm_example_with_layouts<APrecType, BPrecType, CPrecType>(
argc, argv, Col{}, Col{}, Row{});
}
#endif
else
{
throw std::runtime_error("Unsupported memory layout for the input matrices!");
@@ -312,6 +320,7 @@ int run_gemm_example(int argc, char* argv[])
{
return run_gemm_example_prec_type<ck_tile::half_t>(a_layout, b_layout, argc, argv);
}
#if 0
else if(data_type == "bf16")
{
return run_gemm_example_prec_type<ck_tile::bf16_t>(a_layout, b_layout, argc, argv);
@@ -334,6 +343,7 @@ int run_gemm_example(int argc, char* argv[])
return run_gemm_example_prec_type<ck_tile::half_t, ck_tile::pk_int4_t, ck_tile::half_t>(
a_layout, b_layout, argc, argv);
}
#endif
#endif
else
{

6
include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops.hpp
View File

@@ -9,6 +9,7 @@
#include "ck/tensor_operation/gpu/warp/xdlops_gemm.hpp"
#include "ck/tensor_description/tensor_adaptor.hpp"
#define ENABLE_HOT 1
// Double LDS buffer
// Prefetech 2 stage
// Local prefetch 1 stage
@@ -372,6 +373,7 @@ struct BlockwiseGemmXdlops_pipeline_v4
__device__ static constexpr auto HotLoopScheduler()
{
#if ENABLE_HOT
// schedule
constexpr auto num_ds_read_inst =
HotLoopInstList::A_LDS_Read_Inst_Num + HotLoopInstList::B_LDS_Read_Inst_Num;
@@ -398,13 +400,14 @@ struct BlockwiseGemmXdlops_pipeline_v4
__builtin_amdgcn_sched_group_barrier(
0x008, num_mfma_inst / num_buffer_load_inst - 3, 0); // MFMA
});
#endif
}
template <index_t stage>
__device__ static constexpr auto TailScheduler()
{
}
#if ENABLE_HOT
template <>
__device__ constexpr auto TailScheduler<1>()
{
@@ -449,6 +452,7 @@ struct BlockwiseGemmXdlops_pipeline_v4
0x008, num_mfma_inst / num_ds_read_inst, 0); // MFMA
});
}
#endif
static constexpr AMmaTileDesc a_block_desc_m0_m1_m2_k;
static constexpr BMmaTileDesc b_block_desc_n0_n1_n2_k;

2
include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3.hpp
View File

@@ -357,6 +357,7 @@ struct DeviceGemm_Xdl_CShuffleV3 : public DeviceGemmV2<ALayout,
if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v1 ||
BlkGemmPipelineVer == BlockGemmPipelineVersion::v3)
{
#if 0
if(arg.KBatch > 1)
{
const auto kernel =
@@ -367,6 +368,7 @@ struct DeviceGemm_Xdl_CShuffleV3 : public DeviceGemmV2<ALayout,
Run(kernel);
}
else
#endif
{
const auto kernel =
kernel_gemm_xdl_cshuffle_v3<GridwiseGemm,

8
include/ck_tile/core/arch/amd_buffer_addressing.hpp
View File

@@ -38,10 +38,10 @@ CK_TILE_DEVICE int32x4_t make_wave_buffer_resource(const void* ptr, uint32_t siz
{
buffer_resource res{ptr, size, CK_TILE_BUFFER_RESOURCE_3RD_DWORD};
int32x4_t r = __builtin_bit_cast(int32x4_t, res);
r.x = __builtin_amdgcn_readfirstlane(r.x);
r.y = __builtin_amdgcn_readfirstlane(r.y);
r.z = __builtin_amdgcn_readfirstlane(r.z);
r.w = __builtin_amdgcn_readfirstlane(r.w);
//r.x = __builtin_amdgcn_readfirstlane(r.x);
//r.y = __builtin_amdgcn_readfirstlane(r.y);
//r.z = __builtin_amdgcn_readfirstlane(r.z);
//r.w = __builtin_amdgcn_readfirstlane(r.w);
return r;
}

2
include/ck_tile/core/arch/arch.hpp
View File

@@ -70,7 +70,7 @@ CK_TILE_DEVICE index_t get_lane_id() { return __lane_id(); }
CK_TILE_DEVICE index_t get_warp_id()
{
return __builtin_amdgcn_readfirstlane(threadIdx.x / get_warp_size());
return threadIdx.x / get_warp_size();
}
CK_TILE_DEVICE index_t get_thread_id() { return threadIdx.x; }