mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-12 09:16:52 +00:00
061ac0649c
* Few small fixes. * New GroupedGemm instances (BF16) * Unify and refactor GroupedGEMM device API. * Adapt changes to new API. * Adapt grouped gemm profiler. * Accept multiple kbatches for grouped gemm profiler. - delete obsolete two stage as it is now covered by grouped gemm * Update unit test for grouped gemm. * Fix thresholds for BF16 and F8. Unblock tests. * Fix few instances. * Multiple small fixes. * Adapt to new API, check dynamic casting. * Uncomment few data types in grouped gemm profiler. * Fix call to SetDeviceArgs. * Fix profile grouped gemm multiply tile loop. * Fix grouped gemm tile loop kernel args in client examples. * Review comments.
38 lines
804 B
C++
38 lines
804 B
C++
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
|
|
#include <ostream>
|
|
|
|
#pragma once
|
|
|
|
#include "ck/utility/common_header.hpp"
|
|
|
|
namespace ck {
|
|
|
|
enum struct LoopScheduler
|
|
{
|
|
Default,
|
|
Interwave,
|
|
};
|
|
|
|
constexpr LoopScheduler make_default_loop_scheduler()
|
|
{
|
|
#if CK_EXPERIMENTAL_DEFAULT_TO_INTER_WAVE_SCHEDULING
|
|
return LoopScheduler::Interwave;
|
|
#else
|
|
return LoopScheduler::Default;
|
|
#endif // if CK_EXPERIMENTAL_DEFAULT_TO_INTER_WAVE_SCHEDULING
|
|
}
|
|
|
|
} // namespace ck
|
|
|
|
inline std::ostream& operator<<(std::ostream& os, const ck::LoopScheduler& s)
|
|
{
|
|
switch(s)
|
|
{
|
|
case ck::LoopScheduler::Default: os << "Default"; break;
|
|
case ck::LoopScheduler::Interwave: os << "Interwave"; break;
|
|
default: os << "";
|
|
}
|
|
return os;
|
|
}
|