ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-29 19:28:33 +00:00
Code Issues Packages Projects Releases Wiki Activity

commented the 01_add example

Browse Source
This commit is contained in:
AviralGoelAMD
2025-04-15 14:45:34 +00:00
committed by Philip Maybank
parent b1bada8304
commit 1056a980bc
2 changed files with 30 additions and 30 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
example/ck_tile/99_toy_example/01_add/add.cpp
View File

@@ -46,12 +46,12 @@ bool run(const ck_tile::ArgParser& arg_parser)
x_buf_a.ToDevice(x_host_a.data());
x_buf_b.ToDevice(x_host_b.data());
using BlockWarps = ck_tile::sequence<1, 8>;
using BlockTile = ck_tile::sequence<1, 4096>;
using WarpTile = ck_tile::sequence<1, 512>;
using Vector = ck_tile::sequence<1, 8>;
using BlockWarps = ck_tile::sequence<1, 8>; // number of concurrent warps in one block (if 8 warps * 64 threads per warp, 512 threads in one block are NEEDED)
using BlockTile = ck_tile::sequence<1, 4096>; // shape of one blockTile (elements covered by one block)
using WarpTile = ck_tile::sequence<1, 512>; // shape of one warpTile (elements covered by one warp (64 threads))
using Vector = ck_tile::sequence<1, 8>; // shape of one vector (elements covered by one thread)
constexpr ck_tile::index_t kBlockSize = 512;
constexpr ck_tile::index_t kBlockSize = 512; // number of blockWarps * number of threads per warp
constexpr ck_tile::index_t kBlockPerCu = 1;
ck_tile::index_t kGridSize = (m / BlockTile::at(ck_tile::number<0>{}));
std::cout << "block x-size = " << BlockTile::at(ck_tile::number<0>{}) << std::endl;

50
example/ck_tile/99_toy_example/01_add/add.hpp
View File

@@ -14,35 +14,35 @@ template <typename BlockWarps, // num warps along seq<M, N>
typename Vector> // contiguous pixels(vector size) along seq<M, N>
struct AddShape
{
static constexpr index_t Block_M = BlockTile::at(number<0>{});
static constexpr index_t Block_N = BlockTile::at(number<1>{});
static constexpr index_t Block_M = BlockTile::at(number<0>{}); // elements along M in one Block
static constexpr index_t Block_N = BlockTile::at(number<1>{}); // elements along N in one Block
static constexpr index_t Warp_M = WarpTile::at(number<0>{});
static constexpr index_t Warp_N = WarpTile::at(number<1>{});
static constexpr index_t Warp_M = WarpTile::at(number<0>{}); // elements along M in one Warp
static constexpr index_t Warp_N = WarpTile::at(number<1>{}); // elements along N in one Warp
static constexpr index_t Vector_M = Vector::at(number<0>{});
static constexpr index_t Vector_N = Vector::at(number<1>{});
static constexpr index_t Vector_M = Vector::at(number<0>{}); // elements along M in one Vector
static constexpr index_t Vector_N = Vector::at(number<1>{}); // elements along N in one Vector
static constexpr index_t WarpPerBlock_M = BlockWarps::at(number<0>{});
static constexpr index_t WarpPerBlock_N = BlockWarps::at(number<1>{});
static constexpr index_t WarpPerBlock_M = BlockWarps::at(number<0>{}); // num concurrent warps along M
static constexpr index_t WarpPerBlock_N = BlockWarps::at(number<1>{}); // num concurrent warps along N
static constexpr index_t ThreadPerWarp_M = Warp_M / Vector_M;
static constexpr index_t ThreadPerWarp_N = Warp_N / Vector_N;
static constexpr index_t ThreadPerWarp_M = Warp_M / Vector_M; // num threads along M in one Warp (ThreadPerWarp_M * ThreadPerWarp_N must be 64)
static constexpr index_t ThreadPerWarp_N = Warp_N / Vector_N; // num threads along N in one Warp (ThreadPerWarp_M * ThreadPerWarp_N must be 64)
static constexpr index_t Repeat_M = Block_M / (WarpPerBlock_M * Warp_M);
static constexpr index_t Repeat_N = Block_N / (WarpPerBlock_N * Warp_N);
static constexpr index_t Repeat_M = Block_M / (WarpPerBlock_M * Warp_M); // num of time a warp iterates along M to ensure the entire block is covered
static constexpr index_t Repeat_N = Block_N / (WarpPerBlock_N * Warp_N); // num of time a warp iterates along N to ensure the entire block is covered
static constexpr index_t BlockSize =
warpSize * reduce_on_sequence(BlockWarps{}, multiplies{}, number<1>{});
warpSize * reduce_on_sequence(BlockWarps{}, multiplies{}, number<1>{}); // num of threads in one block
};
template <typename XDataType_, typename ComputeDataType_, typename YDataType_, typename BlockShape_>
struct AddProblem
{
using XDataType = remove_cvref_t<XDataType_>;
using ComputeDataType = remove_cvref_t<ComputeDataType_>;
using YDataType = remove_cvref_t<YDataType_>;
using BlockShape = remove_cvref_t<BlockShape_>;
using XDataType = remove_cvref_t<XDataType_>; // data type of input tensor
using ComputeDataType = remove_cvref_t<ComputeDataType_>; // data type of compute tensor
using YDataType = remove_cvref_t<YDataType_>; // data type of output tensor
using BlockShape = remove_cvref_t<BlockShape_>; // block shapes and sizes
};
struct AddDefaultPolicy
@@ -54,12 +54,12 @@ struct AddDefaultPolicy
return make_static_tile_distribution(
tile_distribution_encoding<
sequence<>,
tuple<sequence<S::Repeat_M, S::WarpPerBlock_M, S::ThreadPerWarp_M, S::Vector_M>,
sequence<S::Repeat_N, S::WarpPerBlock_N, S::ThreadPerWarp_N, S::Vector_N>>,
tuple<sequence<1, 2>, sequence<1, 2>>,
tuple<sequence<1, 1>, sequence<2, 2>>,
sequence<1, 1, 2, 2>,
sequence<0, 3, 0, 3>>{});
tuple<sequence<S::Repeat_M, S::WarpPerBlock_M, S::ThreadPerWarp_M, S::Vector_M>, // how many sub division is a block divided in
sequence<S::Repeat_N, S::WarpPerBlock_N, S::ThreadPerWarp_N, S::Vector_N>>, // how many sub division is a block divided in
tuple<sequence<1, 2>, sequence<1, 2>>, // What are the shapes of those sub divisions
tuple<sequence<1, 1>, sequence<2, 2>>, // What are the shapes of those sub divisions
sequence<1, 1, 2, 2>, // How much data does a thread work on and how many iterations of warps are there
sequence<0, 3, 0, 3>>{}); // How much data does a thread work on and how many iterations of warps are there
}
};
@@ -81,7 +81,7 @@ struct Add
const auto x_m_n_a = make_naive_tensor_view<address_space_enum::global,
memory_operation_enum::set,
amd_buffer_coherence_enum::slc>(
p_x_a, make_tuple(M, N), make_tuple(N, 1), number<S::Vector_N>{}, number<1>{});
p_x_a, make_tuple(M, N), make_tuple(N, 1), number<S::Vector_N>{}, number<1>{}); // raw data, shape of tensor, stride of tensor, lastGarunteedVectorLength, lastGarunteedVectorStride
const auto x_m_n_b = make_naive_tensor_view<address_space_enum::global,
memory_operation_enum::set,
@@ -93,7 +93,7 @@ struct Add
amd_buffer_coherence_enum::slc>(
p_y, make_tuple(M, N), make_tuple(N, 1), number<S::Vector_N>{}, number<1>{});
const auto iM = get_block_id() * S::Block_M;
const auto iM = get_block_id() * S::Block_M; // origin of the block along
auto x_window_a = make_tile_window(x_m_n_a,
make_tuple(number<S::Block_M>{}, number<S::Block_N>{}),