mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-03 13:11:25 +00:00
a3499e38b2
Addition of initial CK Tile Stream-K example for bf16 and fp16. These examples are minimal. As more functionality and gtests are added for Stream-K (coming in future PRs), these examples will be expanded.
133 lines
4.6 KiB
C++
133 lines
4.6 KiB
C++
// Copyright © Advanced Micro Devices, Inc., or its affiliates.
|
|
// SPDX-License-Identifier: MIT
|
|
|
|
#pragma once
|
|
|
|
#include <string>
|
|
|
|
#include "ck_tile/core.hpp"
|
|
#include "ck_tile/host/kernel_launch.hpp"
|
|
#include "ck_tile/ops/epilogue.hpp"
|
|
#include "ck_tile/ops/gemm.hpp"
|
|
|
|
#define CK_TILE_PIPELINE_MEMORY 1
|
|
|
|
struct GemmConfigBase
|
|
{
|
|
static constexpr bool kPadM = true;
|
|
static constexpr bool kPadN = true;
|
|
static constexpr bool kPadK = true;
|
|
|
|
static constexpr bool PermuteA = false;
|
|
static constexpr bool PermuteB = false;
|
|
|
|
static constexpr bool TransposeC = false;
|
|
static constexpr bool UseStructuredSparsity = false;
|
|
static constexpr bool Persistent = false;
|
|
|
|
static constexpr int kBlockPerCu = 1;
|
|
static constexpr auto Scheduler = ck_tile::GemmPipelineScheduler::Intrawave;
|
|
static constexpr ck_tile::index_t Pipeline = CK_TILE_PIPELINE_MEMORY;
|
|
static constexpr ck_tile::index_t NumWaveGroups = 1;
|
|
static constexpr bool Preshuffle = false;
|
|
static constexpr bool DoubleSmemBuffer = false;
|
|
};
|
|
|
|
template <typename PrecType>
|
|
struct GemmConfigMemoryInterwave : public GemmConfigBase
|
|
{
|
|
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 = 32;
|
|
|
|
static constexpr ck_tile::index_t M_Warp = 2;
|
|
static constexpr ck_tile::index_t N_Warp = 2;
|
|
static constexpr ck_tile::index_t K_Warp = 1;
|
|
|
|
static constexpr ck_tile::index_t M_Warp_Tile = 32;
|
|
static constexpr ck_tile::index_t N_Warp_Tile = 32;
|
|
static constexpr ck_tile::index_t K_Warp_Tile = sizeof(PrecType) == 2 ? 8 : 16;
|
|
|
|
static constexpr ck_tile::index_t Pipeline = CK_TILE_PIPELINE_MEMORY;
|
|
static constexpr auto Scheduler = ck_tile::GemmPipelineScheduler::Intrawave;
|
|
};
|
|
|
|
template <ck_tile::index_t PipelineId>
|
|
struct PipelineTypeTraits;
|
|
|
|
template <>
|
|
struct PipelineTypeTraits<CK_TILE_PIPELINE_MEMORY>
|
|
{
|
|
template <typename PipelineProblem>
|
|
using GemmPipeline = ck_tile::GemmPipelineAgBgCrMem<PipelineProblem>;
|
|
template <typename PipelineProblem>
|
|
using UniversalGemmPipeline = ck_tile::BaseGemmPipelineAgBgCrMem<PipelineProblem>;
|
|
};
|
|
|
|
template <typename ADataType_, typename BDataType_ = ADataType_, typename CDataType_ = ADataType_>
|
|
struct StreamKGemmTypeConfig
|
|
{
|
|
using ADataType = ADataType_;
|
|
using BDataType = BDataType_;
|
|
using AccDataType = float;
|
|
using CDataType = CDataType_;
|
|
};
|
|
|
|
template <typename T>
|
|
struct DataTypeTraits;
|
|
|
|
template <>
|
|
struct DataTypeTraits<float>
|
|
{
|
|
static constexpr const char* name = "fp32";
|
|
};
|
|
|
|
template <>
|
|
struct DataTypeTraits<ck_tile::half_t>
|
|
{
|
|
static constexpr const char* name = "fp16";
|
|
};
|
|
|
|
template <>
|
|
struct DataTypeTraits<ck_tile::bf16_t>
|
|
{
|
|
static constexpr const char* name = "bf16";
|
|
};
|
|
|
|
auto create_args(int argc, char* argv[])
|
|
{
|
|
ck_tile::ArgParser arg_parser;
|
|
arg_parser.insert("m", "512", "m dimension")
|
|
.insert("n", "512", "n dimension")
|
|
.insert("k", "512", "k dimension")
|
|
.insert("a_layout", "R", "A tensor data layout - Row by default")
|
|
.insert("b_layout", "C", "B tensor data layout - Column by default")
|
|
.insert("c_layout", "R", "C tensor data layout - Row by default")
|
|
.insert("num_sk_blocks",
|
|
"-1",
|
|
"number of Stream-K blocks. -1: chosen by algorithm, or user selected")
|
|
.insert("reduction_strategy",
|
|
"atomic",
|
|
"strategy for storing results in C tensor - atomic/reduction")
|
|
.insert(
|
|
"occupancy",
|
|
"-1",
|
|
"maximum number of workgroups per CU - value of -1 queries occupancy from the device")
|
|
.insert("num_cu",
|
|
"-1",
|
|
"number of compute units (CUs) - value of -1 uses number of CUs on the device")
|
|
.insert("stride_a", "0", "Tensor A stride")
|
|
.insert("stride_b", "0", "Tensor B stride")
|
|
.insert("stride_c", "0", "Tensor C stride")
|
|
.insert("v", "2", "0. No validation, 1. Validation on CPU, 2. Validation on GPU")
|
|
.insert("prec", "fp16", "data type. fp16/bf16")
|
|
.insert("warmup", "50", "number of iterations before benchmarking the kernel")
|
|
.insert("repeat", "100", "number of iterations to benchmark the kernel")
|
|
.insert("timer", "gpu", "gpu:gpu timer, cpu:cpu timer")
|
|
.insert("init", "0", "0:random, 1:linear, 2:constant(1)")
|
|
.insert("flush_cache", "true", "flush cache before running the kernel, defaults to true");
|
|
|
|
bool result = arg_parser.parse(argc, argv);
|
|
return std::make_tuple(result, arg_parser);
|
|
}
|