mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-14 10:09:41 +00:00
* Adding RapidJson Library
* Adding Json Dumps in all CK_Tile Examples
Not verified yet
* Adding json to cktile Batched Transpose
* adding json dumps to layernorm2d_fwd
* Adding json dump to flatmm_basic
* Adding RapidJson Library
* Adding Json Dumps in all CK_Tile Examples
Not verified yet
* Adding json to cktile Batched Transpose
* adding json dumps to layernorm2d_fwd
* Adding json dump to flatmm_basic
* Adding json in 03_gemm
* Add json dump to 16_batched_gemm
* Add json dump to gemm_multi_d_fp16
* Add json dump to grouped_gemm
* fix fmha_bwd/fwd
* Fix clang-format errors
exclude include/rapidjson in jenkins as its a third-party library
* Saparating function and defination.
* Update Documentation of 03_gemm
* Refactoring as per code review
* Disable fp8 instances on unsupported targets (#2592)
* Restrict building of gemm_universal_preshuffle_f8 instances to specific targets in CMakeLists.txt
* Add condition to skip gemm_xdl_universal_preshuffle_f8 instances for unsupported targets in CMakeLists.txt
* Add conditions to skip unsupported targets for gemm_universal_preshuffle_f8 and gemm_xdl_universal_preshuffle_f8 instances in CMakeLists.txt
* Refine conditions to exclude gemm_universal_preshuffle_f8 instances for unsupported targets in CMakeLists.txt
---------
Co-authored-by: AviralGoelAMD <aviralgoel@amd.com>
* fix clang format
* remove duplicate lines of code from library/src/tensor_operation_instance/gpu/CMakeLists.txt
* Fixing Readme and unifying jsondumps
* adding moe_smoothquant
* adding fused_moe
* Fixing Readme for batched_gemm
* Fixing Readme for grouped_gemm
* adding flatmm
* adding gemm_multi_d_fp16
* adding elementwise
* adding File name when json is dumped
* Fixing Reduce after merge
* adding batched_transpose
* Adding Warptile in Gemm
* Fixing Clang Format
---------
Co-authored-by: Aviral Goel <aviral.goel@amd.com>
Co-authored-by: AviralGoelAMD <aviralgoel@amd.com>
Co-authored-by: illsilin_amdeng <Illia.Silin@amd.com>
[ROCm/composable_kernel commit: 4d041837ad]
187 lines
6.5 KiB
C++
187 lines
6.5 KiB
C++
// Tencent is pleased to support the open source community by making RapidJSON available.
|
|
//
|
|
// Copyright (C) 2015 THL A29 Limited, a Tencent company, and Milo Yip.
|
|
//
|
|
// Licensed under the MIT License (the "License"); you may not use this file except
|
|
// in compliance with the License. You may obtain a copy of the License at
|
|
//
|
|
// http://opensource.org/licenses/MIT
|
|
//
|
|
// Unless required by applicable law or agreed to in writing, software distributed
|
|
// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
|
|
// CONDITIONS OF ANY KIND, either express or implied. See the License for the
|
|
// specific language governing permissions and limitations under the License.
|
|
|
|
#ifndef RAPIDJSON_INTERNAL_META_H_
|
|
#define RAPIDJSON_INTERNAL_META_H_
|
|
|
|
#include "../rapidjson.h"
|
|
|
|
#ifdef __GNUC__
|
|
RAPIDJSON_DIAG_PUSH
|
|
RAPIDJSON_DIAG_OFF(effc++)
|
|
#endif
|
|
|
|
#if defined(_MSC_VER) && !defined(__clang__)
|
|
RAPIDJSON_DIAG_PUSH
|
|
RAPIDJSON_DIAG_OFF(6334)
|
|
#endif
|
|
|
|
#if RAPIDJSON_HAS_CXX11_TYPETRAITS
|
|
#include <type_traits>
|
|
#endif
|
|
|
|
//@cond RAPIDJSON_INTERNAL
|
|
RAPIDJSON_NAMESPACE_BEGIN
|
|
namespace internal {
|
|
|
|
// Helper to wrap/convert arbitrary types to void, useful for arbitrary type matching
|
|
template <typename T> struct Void { typedef void Type; };
|
|
|
|
///////////////////////////////////////////////////////////////////////////////
|
|
// BoolType, TrueType, FalseType
|
|
//
|
|
template <bool Cond> struct BoolType {
|
|
static const bool Value = Cond;
|
|
typedef BoolType Type;
|
|
};
|
|
typedef BoolType<true> TrueType;
|
|
typedef BoolType<false> FalseType;
|
|
|
|
|
|
///////////////////////////////////////////////////////////////////////////////
|
|
// SelectIf, BoolExpr, NotExpr, AndExpr, OrExpr
|
|
//
|
|
|
|
template <bool C> struct SelectIfImpl { template <typename T1, typename T2> struct Apply { typedef T1 Type; }; };
|
|
template <> struct SelectIfImpl<false> { template <typename T1, typename T2> struct Apply { typedef T2 Type; }; };
|
|
template <bool C, typename T1, typename T2> struct SelectIfCond : SelectIfImpl<C>::template Apply<T1,T2> {};
|
|
template <typename C, typename T1, typename T2> struct SelectIf : SelectIfCond<C::Value, T1, T2> {};
|
|
|
|
template <bool Cond1, bool Cond2> struct AndExprCond : FalseType {};
|
|
template <> struct AndExprCond<true, true> : TrueType {};
|
|
template <bool Cond1, bool Cond2> struct OrExprCond : TrueType {};
|
|
template <> struct OrExprCond<false, false> : FalseType {};
|
|
|
|
template <typename C> struct BoolExpr : SelectIf<C,TrueType,FalseType>::Type {};
|
|
template <typename C> struct NotExpr : SelectIf<C,FalseType,TrueType>::Type {};
|
|
template <typename C1, typename C2> struct AndExpr : AndExprCond<C1::Value, C2::Value>::Type {};
|
|
template <typename C1, typename C2> struct OrExpr : OrExprCond<C1::Value, C2::Value>::Type {};
|
|
|
|
|
|
///////////////////////////////////////////////////////////////////////////////
|
|
// AddConst, MaybeAddConst, RemoveConst
|
|
template <typename T> struct AddConst { typedef const T Type; };
|
|
template <bool Constify, typename T> struct MaybeAddConst : SelectIfCond<Constify, const T, T> {};
|
|
template <typename T> struct RemoveConst { typedef T Type; };
|
|
template <typename T> struct RemoveConst<const T> { typedef T Type; };
|
|
|
|
|
|
///////////////////////////////////////////////////////////////////////////////
|
|
// IsSame, IsConst, IsMoreConst, IsPointer
|
|
//
|
|
template <typename T, typename U> struct IsSame : FalseType {};
|
|
template <typename T> struct IsSame<T, T> : TrueType {};
|
|
|
|
template <typename T> struct IsConst : FalseType {};
|
|
template <typename T> struct IsConst<const T> : TrueType {};
|
|
|
|
template <typename CT, typename T>
|
|
struct IsMoreConst
|
|
: AndExpr<IsSame<typename RemoveConst<CT>::Type, typename RemoveConst<T>::Type>,
|
|
BoolType<IsConst<CT>::Value >= IsConst<T>::Value> >::Type {};
|
|
|
|
template <typename T> struct IsPointer : FalseType {};
|
|
template <typename T> struct IsPointer<T*> : TrueType {};
|
|
|
|
///////////////////////////////////////////////////////////////////////////////
|
|
// IsBaseOf
|
|
//
|
|
#if RAPIDJSON_HAS_CXX11_TYPETRAITS
|
|
|
|
template <typename B, typename D> struct IsBaseOf
|
|
: BoolType< ::std::is_base_of<B,D>::value> {};
|
|
|
|
#else // simplified version adopted from Boost
|
|
|
|
template<typename B, typename D> struct IsBaseOfImpl {
|
|
RAPIDJSON_STATIC_ASSERT(sizeof(B) != 0);
|
|
RAPIDJSON_STATIC_ASSERT(sizeof(D) != 0);
|
|
|
|
typedef char (&Yes)[1];
|
|
typedef char (&No) [2];
|
|
|
|
template <typename T>
|
|
static Yes Check(const D*, T);
|
|
static No Check(const B*, int);
|
|
|
|
struct Host {
|
|
operator const B*() const;
|
|
operator const D*();
|
|
};
|
|
|
|
enum { Value = (sizeof(Check(Host(), 0)) == sizeof(Yes)) };
|
|
};
|
|
|
|
template <typename B, typename D> struct IsBaseOf
|
|
: OrExpr<IsSame<B, D>, BoolExpr<IsBaseOfImpl<B, D> > >::Type {};
|
|
|
|
#endif // RAPIDJSON_HAS_CXX11_TYPETRAITS
|
|
|
|
|
|
//////////////////////////////////////////////////////////////////////////
|
|
// EnableIf / DisableIf
|
|
//
|
|
template <bool Condition, typename T = void> struct EnableIfCond { typedef T Type; };
|
|
template <typename T> struct EnableIfCond<false, T> { /* empty */ };
|
|
|
|
template <bool Condition, typename T = void> struct DisableIfCond { typedef T Type; };
|
|
template <typename T> struct DisableIfCond<true, T> { /* empty */ };
|
|
|
|
template <typename Condition, typename T = void>
|
|
struct EnableIf : EnableIfCond<Condition::Value, T> {};
|
|
|
|
template <typename Condition, typename T = void>
|
|
struct DisableIf : DisableIfCond<Condition::Value, T> {};
|
|
|
|
// SFINAE helpers
|
|
struct SfinaeTag {};
|
|
template <typename T> struct RemoveSfinaeTag;
|
|
template <typename T> struct RemoveSfinaeTag<SfinaeTag&(*)(T)> { typedef T Type; };
|
|
|
|
#define RAPIDJSON_REMOVEFPTR_(type) \
|
|
typename ::RAPIDJSON_NAMESPACE::internal::RemoveSfinaeTag \
|
|
< ::RAPIDJSON_NAMESPACE::internal::SfinaeTag&(*) type>::Type
|
|
|
|
#define RAPIDJSON_ENABLEIF(cond) \
|
|
typename ::RAPIDJSON_NAMESPACE::internal::EnableIf \
|
|
<RAPIDJSON_REMOVEFPTR_(cond)>::Type * = NULL
|
|
|
|
#define RAPIDJSON_DISABLEIF(cond) \
|
|
typename ::RAPIDJSON_NAMESPACE::internal::DisableIf \
|
|
<RAPIDJSON_REMOVEFPTR_(cond)>::Type * = NULL
|
|
|
|
#define RAPIDJSON_ENABLEIF_RETURN(cond,returntype) \
|
|
typename ::RAPIDJSON_NAMESPACE::internal::EnableIf \
|
|
<RAPIDJSON_REMOVEFPTR_(cond), \
|
|
RAPIDJSON_REMOVEFPTR_(returntype)>::Type
|
|
|
|
#define RAPIDJSON_DISABLEIF_RETURN(cond,returntype) \
|
|
typename ::RAPIDJSON_NAMESPACE::internal::DisableIf \
|
|
<RAPIDJSON_REMOVEFPTR_(cond), \
|
|
RAPIDJSON_REMOVEFPTR_(returntype)>::Type
|
|
|
|
} // namespace internal
|
|
RAPIDJSON_NAMESPACE_END
|
|
//@endcond
|
|
|
|
#if defined(_MSC_VER) && !defined(__clang__)
|
|
RAPIDJSON_DIAG_POP
|
|
#endif
|
|
|
|
#ifdef __GNUC__
|
|
RAPIDJSON_DIAG_POP
|
|
#endif
|
|
|
|
#endif // RAPIDJSON_INTERNAL_META_H_
|