mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-12 17:26:00 +00:00
* updating codegen build for MIOpen access: adding .cmake for codegen component * updating CMake * adding in header guards for some headers due to issues with hiprtc compilation in MIOpen * some more header guards * putting env file in header guard * cleaning up some includes * updated types file for hiprtc purposes * fixed types file: bit-wise/memcpy issue * updating multiple utility files to deal with standard header inclusion for hiprtc * added some more header guards in the utility files, replacing some standard header functionality * added some more header guards * fixing some conflicts in utility files, another round of header guards * fixing errors in data type file * resolved conflict errors in a few utility files * added header guards/replicated functionality in device files * resolved issues with standard headers in device files: device_base and device_grouped_conv_fwd_multiple_abd * resolved issues with standard headers in device files: device_base.hpp, device_grouped_conv_fwd_multiple_abd.hpp, device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp * added header guards for gridwise gemm files: gridwise_gemm_multiple_abd_xdl_cshuffle.hpp and gridwise_gemm_multiple_d_xdl_cshuffle.hpp * fixed issue with numerics header, removed from transform_conv_fwd_to_gemm and added to device_column_to_image_impl, device_grouped_conv_fwd_multiple_abd_xdl_cshuffle, device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3, device_image_to_column_impl * replaced standard header usage and added header guards in block to ctile map and gridwise_gemm_pipeline_selector * resolved errors in device_gemm_xdl_splitk_c_shuffle files in regards to replacement of standard headers in previous commit * added replicated functionality for standard header methods in utility files * replaced standard header functionality in threadwise tensor slice transfer files and added header guards in element_wise_operation.hpp * temp fix for namespace error in MIOpen * remove standard header usage in codegen device op * removed standard header usage in elementwise files, resolved namespace errors * formatting fix * changed codegen argument to ON for testing * temporarily removing codegen compiler flag for testing purposes * added codegen flag again, set default to ON * set codegen flag default back to OFF * replaced enable_if_t standard header usage in data_type.hpp * added some debug prints to pinpoint issues in MIOpen * added print outs to debug in MIOpen * removed debug print outs from device op * resolved stdexcept include error * formatting fix * adding includes to new fp8 file to resolve ck::enable_if_t errors * made changes to amd_wave_read_first_lane * updated functionality in type utility file * fixed end of file issue * resovled errors in type utility file, added functionality to array utility file * fixed standard header usage replication in data_type file, resolves error with failing examples on navi3x * formatting fix * replaced standard header usage in amd_ck_fp8 file * added include to random_gen file * removed and replicated standard header usage from data_type and type_convert files for fp8 changes * replicated standard unsigned integer types in random_gen * resolved comments from review: put calls to reinterpret_cast for size_t in header guards * updated/added copyright headers * removed duplicate header * fixed typo in header guard * updated copyright headers --------- Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
188 lines
5.3 KiB
C++
188 lines
5.3 KiB
C++
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
|
|
|
|
#ifndef CK_CODE_GEN_RTC
|
|
#pragma once
|
|
|
|
#include <cstdlib>
|
|
#include <cstring>
|
|
#include <string>
|
|
#include <string_view>
|
|
|
|
namespace ck {
|
|
namespace internal {
|
|
template <typename T>
|
|
struct ParseEnvVal
|
|
{
|
|
};
|
|
|
|
template <>
|
|
struct ParseEnvVal<bool>
|
|
{
|
|
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<unsigned char>(c));
|
|
}
|
|
}
|
|
|
|
if(value_env_str == "disable" || value_env_str == "disabled" || value_env_str == "0" ||
|
|
value_env_str == "no" || value_env_str == "off" || value_env_str == "false")
|
|
{
|
|
return false;
|
|
}
|
|
else if(value_env_str == "enable" || value_env_str == "enabled" || value_env_str == "1" ||
|
|
value_env_str == "yes" || value_env_str == "on" || value_env_str == "true")
|
|
{
|
|
return true;
|
|
}
|
|
else
|
|
{
|
|
throw std::runtime_error("Invalid value for env variable");
|
|
}
|
|
|
|
return false; // shouldn't reach here
|
|
}
|
|
};
|
|
|
|
// 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<uint64_t>
|
|
{
|
|
static uint64_t parse_env_var_value(const char* vp) { return std::strtoull(vp, nullptr, 0); }
|
|
};
|
|
|
|
template <>
|
|
struct ParseEnvVal<std::string>
|
|
{
|
|
static std::string parse_env_var_value(const char* vp) { return std::string{vp}; }
|
|
};
|
|
|
|
template <typename T>
|
|
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<T>::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_DECLARE_ENV_VAR(name, type, default_val) \
|
|
namespace ck::env { \
|
|
struct name \
|
|
{ \
|
|
static_assert(std::is_same_v<name, ::ck::env::name>, \
|
|
"CK_DECLARE_ENV* must be used in the global namespace"); \
|
|
using value_type = type; \
|
|
static ck::internal::EnvVar<type>& Ref() \
|
|
{ \
|
|
static ck::internal::EnvVar<type> var{#name, default_val}; \
|
|
return var; \
|
|
} \
|
|
}; \
|
|
}
|
|
|
|
#define CK_DECLARE_ENV_VAR_BOOL(name) CK_DECLARE_ENV_VAR(name, bool, false)
|
|
|
|
#define CK_DECLARE_ENV_VAR_UINT64(name) CK_DECLARE_ENV_VAR(name, uint64_t, 0)
|
|
|
|
#define CK_DECLARE_ENV_VAR_STR(name) CK_DECLARE_ENV_VAR(name, std::string, "")
|
|
|
|
#define CK_ENV(name) \
|
|
ck::env::name {}
|
|
|
|
template <class EnvVar>
|
|
inline const std::string& EnvGetString(EnvVar)
|
|
{
|
|
static_assert(std::is_same_v<typename EnvVar::value_type, std::string>);
|
|
return EnvVar::Ref().GetValue();
|
|
}
|
|
|
|
template <class EnvVar>
|
|
inline bool EnvIsEnabled(EnvVar)
|
|
{
|
|
static_assert(std::is_same_v<typename EnvVar::value_type, bool>);
|
|
return !EnvVar::Ref().IsUnset() && EnvVar::Ref().GetValue();
|
|
}
|
|
|
|
template <class EnvVar>
|
|
inline bool EnvIsDisabled(EnvVar)
|
|
{
|
|
static_assert(std::is_same_v<typename EnvVar::value_type, bool>);
|
|
return !EnvVar::Ref().IsUnset() && !EnvVar::Ref().GetValue();
|
|
}
|
|
|
|
template <class EnvVar>
|
|
inline uint64_t EnvValue(EnvVar)
|
|
{
|
|
static_assert(std::is_same_v<typename EnvVar::value_type, uint64_t>);
|
|
return EnvVar::Ref().GetValue();
|
|
}
|
|
|
|
template <class EnvVar>
|
|
inline bool EnvIsUnset(EnvVar)
|
|
{
|
|
return EnvVar::Ref().IsUnset();
|
|
}
|
|
|
|
template <class EnvVar>
|
|
void EnvUnset(EnvVar)
|
|
{
|
|
EnvVar::Ref().Unset();
|
|
}
|
|
|
|
/// updates the cached value of an environment variable
|
|
template <typename EnvVar, typename ValueType>
|
|
void UpdateEnvVar(EnvVar, const ValueType& val)
|
|
{
|
|
static_assert(std::is_same_v<typename EnvVar::value_type, ValueType>);
|
|
EnvVar::Ref().UpdateValue(val);
|
|
}
|
|
|
|
template <typename EnvVar>
|
|
void UpdateEnvVar(EnvVar, const std::string_view& val)
|
|
{
|
|
EnvVar::Ref().UpdateValue(
|
|
ck::internal::ParseEnvVal<typename EnvVar::value_type>::parse_env_var_value(val.data()));
|
|
}
|
|
|
|
} // namespace ck
|
|
#endif
|