ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-17 09:08:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

first commit

Browse Source
This commit is contained in:
Gino Lu
2025-07-30 02:36:06 +00:00
parent 9d4b494f07
commit ab80fd3d7f
5 changed files with 112 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
include/ck_tile/core/utility/functional_with_tuple.hpp
View File

@@ -170,4 +170,20 @@ struct static_uford
}
};
namespace detail {
template <std::size_t Begin, std::size_t... Is, typename Tuple>
auto tuple_slice_impl(Tuple&& t, std::index_sequence<Is...>)
{
return std::forward_as_tuple(std::get<Begin + Is>(std::forward<Tuple>(t))...);
}
}
template <std::size_t Begin, std::size_t End, typename Tuple>
auto tuple_slice(Tuple&& t)
{
return detail::tuple_slice_impl<Begin>(
std::forward<Tuple>(t),
std::make_index_sequence<End - Begin>{});
}
} // namespace ck_tile

56
include/ck_tile/host/kernel_launch.hpp
View File

@@ -5,6 +5,7 @@
#include "ck_tile/core/config.hpp"
#include "ck_tile/core/utility/ignore.hpp"
#include "ck_tile/core/utility/functional_with_tuple.hpp"
#include "ck_tile/host/hip_check_error.hpp"
#include "ck_tile/host/stream_config.hpp"
#include "ck_tile/host/timer.hpp"
@@ -62,6 +63,13 @@ CK_TILE_HOST void launch_and_check(const stream_config& sc, Callables&&... calla
HIP_CHECK_ERROR(hipGetLastError());
}
}
template <typename Callables>
CK_TILE_HOST void launch_and_check_tuple(const stream_config& sc, Callables&& callables)
{
std::apply([&](auto&&... fs) {
launch_and_check(sc, std::forward<decltype(fs)>(fs)...);
}, std::forward<Callables>(callables));
}
// clang-format off
/*
@@ -90,10 +98,14 @@ CK_TILE_HOST void launch_and_check(const stream_config& sc, Callables&&... calla
* ...);
**/
// clang-format on
template <typename... Callables>
template <size_t TimerStart, size_t TimerEnd, typename... Callables>
CK_TILE_HOST float launch_kernel(const stream_config& s, Callables&&... callables)
{
static_assert(sizeof...(callables) > 0, "At least one callable is required!");
constexpr auto N_callables = sizeof...(callables);
static_assert(N_callables > 0, "At least one callable is required!");
static_assert(N_callables >= TimerEnd, "Wrong timer range.");
static_assert(TimerEnd >= TimerStart, "Wrong timer range.");
if(!s.time_kernel_)
{
@@ -102,20 +114,41 @@ CK_TILE_HOST float launch_kernel(const stream_config& s, Callables&&... callable
}
auto time_launches = [&](auto timer) {
auto callables_all = std::forward_as_tuple(std::forward<Callables>(callables)...);
auto callables_start = tuple_slice<0, TimerStart>(callables_all);
auto callables_timed = tuple_slice<TimerStart, TimerEnd>(callables_all);
auto callables_end = tuple_slice<TimerEnd, N_callables>(callables_all);
// Warmup
for(int i = 0; i < s.cold_niters_; i++)
{
launch_and_check(s, std::forward<Callables>(callables)...);
launch_and_check_tuple(s, callables_all);
}
timer.start(s.stream_id_);
float times = 0.f;
for(int i = 0; i < s.nrepeat_; i++)
{
launch_and_check(s, std::forward<Callables>(callables)...);
}
timer.stop(s.stream_id_);
if constexpr (std::tuple_size_v<decltype(callables_start)> > 0)
{
launch_and_check_tuple(s, callables_start);
}
return timer.duration() / s.nrepeat_;
if constexpr (std::tuple_size_v<decltype(callables_timed)> > 0)
{
timer.start(s.stream_id_);
launch_and_check_tuple(s, callables_timed);
timer.stop(s.stream_id_);
times += timer.duration();
}
if constexpr (std::tuple_size_v<decltype(callables_end)> > 0)
{
launch_and_check_tuple(s, callables_end);
}
}
return times / s.nrepeat_;
};
if(s.is_gpu_timer_)
@@ -128,6 +161,13 @@ CK_TILE_HOST float launch_kernel(const stream_config& s, Callables&&... callable
}
}
template <size_t TimerStart = 0, typename... Callables>
CK_TILE_HOST float launch_kernel(const stream_config& s, Callables&&... callables)
{
return launch_kernel<TimerStart, sizeof...(Callables)>(
s, std::forward<Callables>(callables)...);
}
template <typename PreprocessFunc, typename... Callables>
CK_TILE_HOST float launch_kernel_preprocess(const stream_config& s,
PreprocessFunc preprocess,

1
test/ck_tile/CMakeLists.txt
View File

@@ -21,3 +21,4 @@ add_subdirectory(add_rmsnorm2d_rdquant)
# add_subdirectory(layernorm2d)
# add_subdirectory(rmsnorm2d)
add_subdirectory(gemm_block_scale)
add_subdirectory(kernel_launch)

1
test/ck_tile/kernel_launch/CMakeLists.txt Normal file
View File

@@ -0,0 +1 @@
add_gtest_executable(test_kernel_launch test_kernel_launch.cpp)

46
test/ck_tile/kernel_launch/test_kernel_launch.cpp Normal file
View File

@@ -0,0 +1,46 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include <hip/hip_runtime.h>
#include "ck_tile/host.hpp"
using ck_tile::stream_config;
TEST(Kernel, old)
{
stream_config s1;
s1.time_kernel_ = true;
s1.cold_niters_ = 1;
s1.nrepeat_ = 2;
ck_tile::launch_kernel(s1, // timer for all kernel
[=](const stream_config& s[[maybe_unused]]) { printf("0, "); },
[=](const stream_config& s[[maybe_unused]]) { printf("1, "); },
[=](const stream_config& s[[maybe_unused]]) { printf("2, "); }
);
}
TEST(Kernel, new_1)
{
stream_config s1;
s1.time_kernel_ = true;
s1.cold_niters_ = 2;
s1.nrepeat_ = 2;
ck_tile::launch_kernel<1, 2>(s1, // timer for kernel [1, 2)
[=](const stream_config& s[[maybe_unused]]) { printf("0, "); },
[=](const stream_config& s[[maybe_unused]]) { printf("1, "); },
[=](const stream_config& s[[maybe_unused]]) { printf("2, "); }
);
}
TEST(Kernel, new_2)
{
stream_config s1;
s1.time_kernel_ = true;
s1.cold_niters_ = 2;
s1.nrepeat_ = 2;
ck_tile::launch_kernel<1>(s1, // timer for kernel [1, N)
[=](const stream_config& s[[maybe_unused]]) { printf("0, "); },
[=](const stream_config& s[[maybe_unused]]) { printf("1, "); },
[=](const stream_config& s[[maybe_unused]]) { printf("2, "); }
);
}