From 45f410af7d776ed468a4563da361a7d2c780a834 Mon Sep 17 00:00:00 2001 From: aledudek Date: Tue, 25 Feb 2025 11:48:38 +0100 Subject: [PATCH] [CK_TILE] Add EnvLogging and missing gemm args checks (#1896) * [CK_TILE] Add EnvLogging - refactor IsSupported error messages * [CK_TILE] Add EnvLogging - wrap gemm kernel error messages * [CK_TILE] Add EnvLogging - Add missing k_batch args check * [CK_TILE] Add EnvLogging - remove debug log * Add one check * [CK_TILE] EnvLogging - add CK_TILE_ERROR logs * [CK_TILE] EnvLogging quotes fix * [CK_TILE] EngLogging use function instead of macro for err logs * [CK_TILE] EnvLogging - refactor checking env var [ROCm/composable_kernel commit: c9bcfd755ed4d2102d76a6f545ac6e9a030d7d8e] --- include/ck_tile/core.hpp | 1 + include/ck_tile/core/config.hpp | 6 + include/ck_tile/core/utility/env.hpp | 204 ++++++++++++++++++ .../ck_tile/ops/gemm/kernel/gemm_kernel.hpp | 89 +++++--- 4 files changed, 273 insertions(+), 27 deletions(-) create mode 100644 include/ck_tile/core/utility/env.hpp diff --git a/include/ck_tile/core.hpp b/include/ck_tile/core.hpp index a8c95b9c38..25f600d68d 100644 --- a/include/ck_tile/core.hpp +++ b/include/ck_tile/core.hpp @@ -58,6 +58,7 @@ #include "ck_tile/core/tensor/transpose_tile.hpp" #include "ck_tile/core/tensor/update_tile.hpp" #include "ck_tile/core/utility/bit_cast.hpp" +#include "ck_tile/core/utility/env.hpp" #include "ck_tile/core/utility/functional.hpp" #include "ck_tile/core/utility/functional_with_tuple.hpp" #include "ck_tile/core/utility/ignore.hpp" diff --git a/include/ck_tile/core/config.hpp b/include/ck_tile/core/config.hpp index c761fcb8c3..090b2bf797 100644 --- a/include/ck_tile/core/config.hpp +++ b/include/ck_tile/core/config.hpp @@ -29,6 +29,12 @@ #include "hip/hip_fp16.h" #endif +#include "ck_tile/core/utility/env.hpp" + +// environment variable to enable logging: +// export CK_TILE_LOGGING=ON or CK_TILE_LOGGING=1 or CK_TILE_LOGGING=ENABLED +CK_TILE_DECLARE_ENV_VAR_BOOL(CK_TILE_LOGGING) + #ifdef __HIPCC__ #define CK_TILE_HOST inline __host__ #define CK_TILE_DEVICE inline __device__ diff --git a/include/ck_tile/core/utility/env.hpp b/include/ck_tile/core/utility/env.hpp new file mode 100644 index 0000000000..5b0b7a9071 --- /dev/null +++ b/include/ck_tile/core/utility/env.hpp @@ -0,0 +1,204 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include +#include + +namespace ck_tile { + +template +void CK_TILE_ERROR(Args&&... args) noexcept +{ + std::ostringstream oss; + (oss << ... << args); + std::cerr << "[ERROR] " << oss.str() << std::endl; +} + +namespace internal { + +template +bool is_any_of(const char* const (&names)[N], const std::string& str) +{ + return std::any_of(std::begin(names), std::end(names), [&](const char* inner_str) { + return str == inner_str; + }); +}; + +template +struct ParseEnvVal +{ +}; +template <> +struct ParseEnvVal +{ + static bool parse_env_var_value(const char* vp) + { + std::string value_env_str{vp}; + + for(auto& c : value_env_str) + { + if(std::isalpha(c) != 0) + { + c = std::tolower(static_cast(c)); + } + } + + if(is_any_of(enabled_names, value_env_str)) + { + return true; + } + else if(is_any_of(disabled_names, value_env_str)) + { + return false; + } + else + { + throw std::runtime_error("Invalid value for env variable"); + } + + return false; + } + + private: + static constexpr const char* enabled_names[] = {"enable", "enabled", "1", "yes", "on", "true"}; + static constexpr const char* disabled_names[] = { + "disable", "disabled", "0", "no", "off", "false"}; +}; + +// Supports hexadecimals (with leading "0x"), octals (if prefix is "0") and decimals (default). +// Returns 0 if environment variable is in wrong format (strtoull fails to parse the string). +template <> +struct ParseEnvVal +{ + static uint64_t parse_env_var_value(const char* vp) { return std::strtoull(vp, nullptr, 0); } +}; + +template <> +struct ParseEnvVal +{ + static std::string parse_env_var_value(const char* vp) { return std::string{vp}; } +}; + +template +struct EnvVar +{ + private: + T value{}; + bool is_unset = true; + + public: + const T& GetValue() const { return value; } + + bool IsUnset() const { return is_unset; } + + void Unset() { is_unset = true; } + + void UpdateValue(const T& val) + { + is_unset = false; + value = val; + } + + explicit EnvVar(const char* const name, const T& def_val) + { + // NOLINTNEXTLINE (concurrency-mt-unsafe) + const char* vp = std::getenv(name); + if(vp != nullptr) // a value was provided + { + is_unset = false; + value = ParseEnvVal::parse_env_var_value(vp); + } + else // no value provided, use default value + { + value = def_val; + } + } +}; +} // end namespace internal + +// Static inside function hides the variable and provides +// thread-safety/locking +// Used in global namespace +#define CK_TILE_DECLARE_ENV_VAR(name, type, default_val) \ + namespace ck_tile::env { \ + struct name \ + { \ + static_assert(std::is_same_v, \ + "CK_TILE_DECLARE_ENV* must be used in the global namespace"); \ + using value_type = type; \ + static ck_tile::internal::EnvVar& Ref() \ + { \ + static ck_tile::internal::EnvVar var{#name, default_val}; \ + return var; \ + } \ + }; \ + } + +#define CK_TILE_DECLARE_ENV_VAR_BOOL(name) CK_TILE_DECLARE_ENV_VAR(name, bool, false) + +#define CK_TILE_DECLARE_ENV_VAR_UINT64(name) CK_TILE_DECLARE_ENV_VAR(name, uint64_t, 0) + +#define CK_TILE_DECLARE_ENV_VAR_STR(name) CK_TILE_DECLARE_ENV_VAR(name, std::string, "") + +#define CK_TILE_ENV(name) \ + ck_tile::env::name {} + +template +inline const std::string& EnvGetString(EnvVar) +{ + static_assert(std::is_same_v); + return EnvVar::Ref().GetValue(); +} + +template +inline bool EnvIsEnabled(EnvVar) +{ + static_assert(std::is_same_v); + return !EnvVar::Ref().IsUnset() && EnvVar::Ref().GetValue(); +} + +template +inline bool EnvIsDisabled(EnvVar) +{ + static_assert(std::is_same_v); + return !EnvVar::Ref().IsUnset() && !EnvVar::Ref().GetValue(); +} + +template +inline uint64_t EnvValue(EnvVar) +{ + static_assert(std::is_same_v); + return EnvVar::Ref().GetValue(); +} + +template +inline bool EnvIsUnset(EnvVar) +{ + return EnvVar::Ref().IsUnset(); +} + +template +void EnvUnset(EnvVar) +{ + EnvVar::Ref().Unset(); +} + +/// Updates the cached value of an environment variable +template +void UpdateEnvVar(EnvVar, const ValueType& val) +{ + static_assert(std::is_same_v); + EnvVar::Ref().UpdateValue(val); +} + +template +void UpdateEnvVar(EnvVar, const std::string_view& val) +{ + EnvVar::Ref().UpdateValue( + ck_tile::internal::ParseEnvVal::parse_env_var_value( + val.data())); +} + +} // namespace ck_tile diff --git a/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp b/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp index 3107d07bc9..741a6b9fc3 100644 --- a/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp +++ b/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp @@ -172,23 +172,32 @@ struct GemmKernel { if(kargs.k_batch != 1) { - std::cerr << "Conditions not met for Kbatch >1 !" << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR("Conditions not met for Kbatch >1 !"); + } return false; } } if constexpr(std::is_same_v) { - if(kargs.K % TilePartitioner::KPerBlock != 0 && GemmPipeline::kPadK == false) + if(kargs.K % (TilePartitioner::KPerBlock * kargs.k_batch) != 0 && + GemmPipeline::kPadK == false) { - std::cerr << "Can't support K that is not a multiple of KPerBlock" - " without padding!" - << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR("Can't support K that is not a multiple of k_batch * KPerBlock " + "without padding!"); + } return false; } if(kargs.K % GemmPipeline::GetVectorSizeA() != 0) { - std::cerr << "K is not a multiple of vector load size for A tensor!" << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR("K is not a multiple of vector load size for A tensor!"); + } return false; } } @@ -196,14 +205,19 @@ struct GemmKernel { if(kargs.M % TilePartitioner::MPerBlock != 0 && GemmPipeline::kPadM == false) { - std::cerr << "Can't support M that is not a multiple of MPerBlock" - " without padding!" - << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR( + "Can't support M that is not a multiple of MPerBlock without padding!"); + } return false; } if(kargs.M % GemmPipeline::GetVectorSizeA() != 0) { - std::cerr << "M is not a multiple of vector load size for A tensor!" << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR("M is not a multiple of vector load size for A tensor!"); + } return false; } } @@ -212,29 +226,40 @@ struct GemmKernel { if(kargs.N % TilePartitioner::NPerBlock != 0 && GemmPipeline::kPadN == false) { - std::cerr << "Can't support N that is not a multiple of NPerBlock" - " without padding!" - << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR( + "Can't support N that is not a multiple of NPerBlock without padding!"); + } return false; } if(kargs.N % GemmPipeline::GetVectorSizeB() != 0) { - std::cerr << "N is not a multiple of vector load size for B tensor!" << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR("N is not a multiple of vector load size for B tensor!"); + } return false; } } else { - if(kargs.K % TilePartitioner::KPerBlock != 0 && GemmPipeline::kPadK == false) + if(kargs.K % (TilePartitioner::KPerBlock * kargs.k_batch) != 0 && + GemmPipeline::kPadK == false) { - std::cerr << "Can't support K that is not a multiple of KPerBlock" - " without padding!" - << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR("Can't support K that is not a multiple of k_batch * KPerBlock " + "without padding!"); + } return false; } if(kargs.K % GemmPipeline::GetVectorSizeB() != 0) { - std::cerr << "K is not a multiple of vector load size for B tensor!" << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR("K is not a multiple of vector load size for B tensor!"); + } return false; } } @@ -243,14 +268,19 @@ struct GemmKernel { if(kargs.N % TilePartitioner::NPerBlock != 0 && GemmPipeline::kPadN == false) { - std::cerr << "Can't support N that is not a multiple of NPerBlock" - " without padding!" - << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR( + "Can't support N that is not a multiple of NPerBlock without padding!"); + } return false; } if(kargs.N % EpiloguePipeline::template GetVectorSizeC() != 0) { - std::cerr << "N is not a multiple of vector load size for C tensor!" << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR("N is not a multiple of vector load size for C tensor!"); + } return false; } } @@ -258,14 +288,19 @@ struct GemmKernel { if(kargs.M % TilePartitioner::MPerBlock != 0 && GemmPipeline::kPadM == false) { - std::cerr << "Can't support M that is not a multiple of MPerBlock" - " without padding!" - << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR( + "Can't support M that is not a multiple of MPerBlock without padding!"); + } return false; } if(kargs.M % EpiloguePipeline::template GetVectorSizeC() != 0) { - std::cerr << "M is not a multiple of vector load size for C tensor!" << std::endl; + if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING))) + { + CK_TILE_ERROR("M is not a multiple of vector load size for C tensor!"); + } return false; } }